From 540ae5967a1a3ec0536b86b6a85139a98d5abeab Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Fri, 7 Jul 2023 19:36:42 +0200 Subject: [PATCH 1/3] Improve dpnp.cos() and dpnp.sin() implementations --- .github/workflows/conda-package.yml | 1 + dpnp/backend/extensions/vm/cos.hpp | 78 +++++++++ dpnp/backend/extensions/vm/sin.hpp | 78 +++++++++ dpnp/backend/extensions/vm/types_matrix.hpp | 38 +++++ dpnp/backend/extensions/vm/vm_py.cpp | 60 +++++++ dpnp/backend/include/dpnp_iface_fptr.hpp | 6 +- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 18 -- dpnp/dpnp_algo/dpnp_algo.pxd | 6 - dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi | 10 -- dpnp/dpnp_algo/dpnp_elementwise_common.py | 100 ++++++++++++ dpnp/dpnp_iface_trigonometric.py | 99 +++++++---- tests/skipped_tests.tbl | 4 + tests/skipped_tests_gpu.tbl | 4 + tests/test_umath.py | 154 ++++++++++++++---- .../cupy/math_tests/test_trigonometric.py | 29 +++- 15 files changed, 580 insertions(+), 105 deletions(-) create mode 100644 dpnp/backend/extensions/vm/cos.hpp create mode 100644 dpnp/backend/extensions/vm/sin.hpp diff --git a/.github/workflows/conda-package.yml b/.github/workflows/conda-package.yml index d2263370280..4470571dc56 100644 --- a/.github/workflows/conda-package.yml +++ b/.github/workflows/conda-package.yml @@ -23,6 +23,7 @@ env: test_umath.py test_usm_type.py third_party/cupy/math_tests/test_explog.py + third_party/cupy/math_tests/test_trigonometric.py third_party/cupy/sorting_tests/test_sort.py VER_JSON_NAME: 'version.json' VER_SCRIPT1: "import json; f = open('version.json', 'r'); j = json.load(f); f.close(); " diff --git a/dpnp/backend/extensions/vm/cos.hpp b/dpnp/backend/extensions/vm/cos.hpp new file mode 100644 index 00000000000..d7212b67e1c --- /dev/null +++ b/dpnp/backend/extensions/vm/cos.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 cos_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::cos(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 CosContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::CosOutputType::value_type, void>) + { + return nullptr; + } + else { + return cos_contig_impl; + } + } +}; +} // namespace vm +} // namespace ext +} // namespace backend +} // namespace dpnp diff --git a/dpnp/backend/extensions/vm/sin.hpp b/dpnp/backend/extensions/vm/sin.hpp new file mode 100644 index 00000000000..90fa3ad114e --- /dev/null +++ b/dpnp/backend/extensions/vm/sin.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 sin_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::sin(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 SinContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename types::SinOutputType::value_type, void>) + { + return nullptr; + } + else { + return sin_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 26235b83872..d7c9f9eecdf 100644 --- a/dpnp/backend/extensions/vm/types_matrix.hpp +++ b/dpnp/backend/extensions/vm/types_matrix.hpp @@ -68,6 +68,25 @@ struct DivOutputType 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::cos function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct CosOutputType +{ + 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; +}; + /** * @brief A factory to define pairs of supported types for which * MKL VM library provides support in oneapi::mkl::vm::ln function. @@ -86,6 +105,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::sin function. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct SinOutputType +{ + 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 59e6612b9ed..0de88cb8b41 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -31,8 +31,10 @@ #include #include "common.hpp" +#include "cos.hpp" #include "div.hpp" #include "ln.hpp" +#include "sin.hpp" #include "types_matrix.hpp" namespace py = pybind11; @@ -43,7 +45,9 @@ 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 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]; PYBIND11_MODULE(_vm_impl, m) { @@ -80,6 +84,34 @@ PYBIND11_MODULE(_vm_impl, m) py::arg("dst")); } + // UnaryUfunc: ==== Cos(x) ==== + { + vm_ext::init_ufunc_dispatch_vector( + cos_dispatch_vector); + + auto cos_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + cos_dispatch_vector); + }; + m.def("_cos", cos_pyapi, + "Call `cos` function from OneMKL VM library to compute " + "natural logarithm of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto cos_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + cos_dispatch_vector); + }; + m.def("_mkl_cos_to_call", cos_need_to_call_pyapi, + "Check input arguments to answer if `cos` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } + // UnaryUfunc: ==== Ln(x) ==== { vm_ext::init_ufunc_dispatch_vector( + sin_dispatch_vector); + + auto sin_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst, + const event_vecT &depends = {}) { + return vm_ext::unary_ufunc(exec_q, src, dst, depends, + sin_dispatch_vector); + }; + m.def("_sin", sin_pyapi, + "Call `sin` function from OneMKL VM library to compute " + "natural logarithm of vector elements", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), + py::arg("depends") = py::list()); + + auto sin_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src, + arrayT dst) { + return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst, + sin_dispatch_vector); + }; + m.def("_mkl_sin_to_call", sin_need_to_call_pyapi, + "Check input arguments to answer if `sin` function from " + "OneMKL VM library can be used", + py::arg("sycl_queue"), py::arg("src"), py::arg("dst")); + } } diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index 335c34bad58..6dec7d944cf 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -145,8 +145,7 @@ enum class DPNPFuncName : size_t DPNP_FN_CORRELATE_EXT, /**< Used in numpy.correlate() impl, requires extra parameters */ DPNP_FN_COS, /**< Used in numpy.cos() impl */ - DPNP_FN_COS_EXT, /**< Used in numpy.cos() impl, requires extra parameters */ - DPNP_FN_COSH, /**< Used in numpy.cosh() impl */ + DPNP_FN_COSH, /**< Used in numpy.cosh() impl */ DPNP_FN_COSH_EXT, /**< Used in numpy.cosh() impl, requires extra parameters */ DPNP_FN_COUNT_NONZERO, /**< Used in numpy.count_nonzero() impl */ @@ -475,8 +474,7 @@ enum class DPNPFuncName : size_t DPNP_FN_SIGN_EXT, /**< Used in numpy.sign() impl, requires extra parameters */ DPNP_FN_SIN, /**< Used in numpy.sin() impl */ - DPNP_FN_SIN_EXT, /**< Used in numpy.sin() impl, requires extra parameters */ - DPNP_FN_SINH, /**< Used in numpy.sinh() impl */ + DPNP_FN_SINH, /**< Used in numpy.sinh() impl */ DPNP_FN_SINH_EXT, /**< Used in numpy.sinh() impl, requires extra parameters */ DPNP_FN_SORT, /**< Used in numpy.sort() impl */ diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 74141e4bd5f..7f6b66afc58 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -486,15 +486,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_COS][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_cos_c_default}; - fmap[DPNPFuncName::DPNP_FN_COS_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_cos_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COS_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_cos_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COS_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_cos_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COS_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_cos_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COSH][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_cosh_c_default}; fmap[DPNPFuncName::DPNP_FN_COSH][eft_LNG][eft_LNG] = { @@ -711,15 +702,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_SIN][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_sin_c_default}; - fmap[DPNPFuncName::DPNP_FN_SIN_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_sin_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SIN_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_sin_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SIN_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_sin_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SIN_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_sin_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SINH][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_sinh_c_default}; fmap[DPNPFuncName::DPNP_FN_SINH][eft_LNG][eft_LNG] = { diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index b04f01f989e..433034f04c6 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -90,8 +90,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_COPYTO_EXT DPNP_FN_CORRELATE DPNP_FN_CORRELATE_EXT - DPNP_FN_COS - DPNP_FN_COS_EXT DPNP_FN_COSH DPNP_FN_COSH_EXT DPNP_FN_COUNT_NONZERO @@ -293,8 +291,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_SEARCHSORTED_EXT DPNP_FN_SIGN DPNP_FN_SIGN_EXT - DPNP_FN_SIN - DPNP_FN_SIN_EXT DPNP_FN_SINH DPNP_FN_SINH_EXT DPNP_FN_SORT @@ -546,7 +542,6 @@ cpdef dpnp_descriptor dpnp_arcsinh(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_arctan(dpnp_descriptor array1, dpnp_descriptor out) cpdef dpnp_descriptor dpnp_arctanh(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_cbrt(dpnp_descriptor array1) -cpdef dpnp_descriptor dpnp_cos(dpnp_descriptor array1, dpnp_descriptor out) cpdef dpnp_descriptor dpnp_cosh(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_degrees(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_exp(dpnp_descriptor array1, dpnp_descriptor out) @@ -557,7 +552,6 @@ cpdef dpnp_descriptor dpnp_log1p(dpnp_descriptor array1) cpdef dpnp_descriptor dpnp_log2(dpnp_descriptor array1) 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) diff --git a/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi b/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi index e8e2e5b16a5..8f4ea3bc80c 100644 --- a/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi +++ b/dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi @@ -43,7 +43,6 @@ __all__ += [ 'dpnp_arctan', 'dpnp_arctanh', 'dpnp_cbrt', - 'dpnp_cos', 'dpnp_cosh', 'dpnp_degrees', 'dpnp_exp', @@ -54,7 +53,6 @@ __all__ += [ 'dpnp_log2', 'dpnp_radians', 'dpnp_recip', - 'dpnp_sin', 'dpnp_sinh', 'dpnp_sqrt', 'dpnp_square', @@ -92,10 +90,6 @@ cpdef utils.dpnp_descriptor dpnp_cbrt(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_CBRT_EXT, x1) -cpdef utils.dpnp_descriptor dpnp_cos(utils.dpnp_descriptor x1, utils.dpnp_descriptor out): - return call_fptr_1in_1out_strides(DPNP_FN_COS_EXT, x1, dtype=None, out=out, where=True, func_name='cos') - - cpdef utils.dpnp_descriptor dpnp_cosh(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_COSH_EXT, x1) @@ -136,10 +130,6 @@ cpdef utils.dpnp_descriptor dpnp_radians(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_RADIANS_EXT, x1) -cpdef utils.dpnp_descriptor dpnp_sin(utils.dpnp_descriptor x1, utils.dpnp_descriptor out): - return call_fptr_1in_1out_strides(DPNP_FN_SIN_EXT, x1, dtype=None, out=out, where=True, func_name='sin') - - cpdef utils.dpnp_descriptor dpnp_sinh(utils.dpnp_descriptor x1): return call_fptr_1in_1out_strides(DPNP_FN_SINH_EXT, x1) diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index 8465454b7ae..5a2f8905d77 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -43,6 +43,7 @@ __all__ = [ "check_nd_call_func", "dpnp_add", + "dpnp_cos", "dpnp_divide", "dpnp_equal", "dpnp_greater", @@ -52,6 +53,7 @@ "dpnp_log", "dpnp_multiply", "dpnp_not_equal", + "dpnp_sin", "dpnp_subtract", ] @@ -167,6 +169,55 @@ def dpnp_add(x1, x2, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_cos_docstring = """ +cos(x, out=None, order='K') +Computes cosine for each element `x_i` for input array `x`. +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + 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 cosine. The data type + of the returned array is determined by the Type Promotion Rules. +""" + + +def dpnp_cos(x, out=None, order="K"): + """ + Invokes cos() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for cos() function. + + """ + + def _call_cos(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_cos_to_call(sycl_queue, src, dst): + # call pybind11 extension for cos() function from OneMKL VM + return vmi._cos(sycl_queue, src, dst, depends) + return ti._cos(src, dst, sycl_queue, depends) + + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + func = UnaryElementwiseFunc( + "cos", ti._cos_result_type, _call_cos, _cos_docstring + ) + res_usm = func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _divide_docstring_ = """ divide(x1, x2, out=None, order="K") @@ -585,6 +636,55 @@ def dpnp_not_equal(x1, x2, out=None, order="K"): return dpnp_array._create_from_usm_ndarray(res_usm) +_sin_docstring = """ +sin(x, out=None, order='K') +Computes sine for each element `x_i` of input array `x`. +Args: + x (dpnp.ndarray): + Input array, expected to have numeric data type. + 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 sine. The data type of the + returned array is determined by the Type Promotion Rules. +""" + + +def dpnp_sin(x, out=None, order="K"): + """ + Invokes sin() function from pybind11 extension of OneMKL VM if possible. + + Otherwise fully relies on dpctl.tensor implementation for sin() function. + + """ + + def _call_sin(src, dst, sycl_queue, depends=None): + """A callback to register in UnaryElementwiseFunc class of dpctl.tensor""" + + if depends is None: + depends = [] + + if vmi._mkl_sin_to_call(sycl_queue, src, dst): + # call pybind11 extension for sin() function from OneMKL VM + return vmi._sin(sycl_queue, src, dst, depends) + return ti._sin(src, dst, sycl_queue, depends) + + # dpctl.tensor only works with usm_ndarray + x1_usm = dpnp.get_usm_ndarray(x) + out_usm = None if out is None else dpnp.get_usm_ndarray(out) + + func = UnaryElementwiseFunc( + "sin", ti._sin_result_type, _call_sin, _sin_docstring + ) + res_usm = func(x1_usm, out=out_usm, order=order) + return dpnp_array._create_from_usm_ndarray(res_usm) + + _subtract_docstring_ = """ subtract(x1, x2, out=None, order="K") diff --git a/dpnp/dpnp_iface_trigonometric.py b/dpnp/dpnp_iface_trigonometric.py index a2be4327fa1..89f6d7c05e0 100644 --- a/dpnp/dpnp_iface_trigonometric.py +++ b/dpnp/dpnp_iface_trigonometric.py @@ -49,7 +49,9 @@ from .dpnp_algo.dpnp_elementwise_common import ( check_nd_call_func, + dpnp_cos, dpnp_log, + dpnp_sin, ) __all__ = [ @@ -408,39 +410,54 @@ def arctan2(x1, x2, dtype=None, out=None, where=True, **kwargs): ) -def cos(x1, out=None, **kwargs): +def cos( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Trigonometric cosine, element-wise. For full documentation refer to :obj:`numpy.cos`. + Returns + ------- + y : dpnp.ndarray + The cosine of each element of `x`. + Limitations ----------- - Input array is supported as :obj:`dpnp.ndarray`. + Parameters `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. Examples -------- >>> import dpnp as np >>> x = np.array([0, np.pi/2, np.pi]) - >>> out = np.cos(x) - >>> [i for i in out] - [1.0, 6.123233995736766e-17, -1.0] + >>> np.cos(x) + array([ 1.000000e+00, -4.371139e-08, -1.000000e+00]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.cos, + dpnp_cos, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) - return dpnp_cos(x1_desc, out_desc).get_pyobj() - - return call_origin(numpy.cos, x1, out=out, **kwargs) def cosh(x1): @@ -945,18 +962,32 @@ def radians(x1): return call_origin(numpy.radians, x1, **kwargs) -def sin(x1, out=None, **kwargs): +def sin( + x, + /, + out=None, + *, + order="K", + where=True, + dtype=None, + subok=True, + **kwargs, +): """ Trigonometric sine, element-wise. For full documentation refer to :obj:`numpy.sin`. + Returns + ------- + y : dpnp.ndarray + The sine of each element of `x`. + Limitations ----------- - Parameters ``x1`` is supported as :obj:`dpnp.ndarray`. - Parameter ``out`` is supported as default value ``None``. - Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. + Parameters `x` is only supported as either :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray`. + Parameters `where`, `dtype` and `subok` are supported with their default values. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. See Also @@ -969,24 +1000,22 @@ def sin(x1, out=None, **kwargs): -------- >>> import dpnp as np >>> x = np.array([0, np.pi/2, np.pi]) - >>> out = np.sin(x) - >>> [i for i in out] - [0.0, 1.0, 1.2246467991473532e-16] + >>> np.sin(x) + array([ 0.000000e+00, 1.000000e+00, -8.742278e-08]) """ - x1_desc = dpnp.get_dpnp_descriptor( - x1, copy_when_strides=False, copy_when_nondefault_queue=False + return check_nd_call_func( + numpy.sin, + dpnp_sin, + x, + out=out, + where=where, + order=order, + dtype=dtype, + subok=subok, + **kwargs, ) - if x1_desc: - out_desc = ( - dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) - if out is not None - else None - ) - return dpnp_sin(x1_desc, out_desc).get_pyobj() - - return call_origin(numpy.sin, x1, out=out, **kwargs) def sinh(x1): diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl index 5fc69258e8b..6f40aa738ec 100644 --- a/tests/skipped_tests.tbl +++ b/tests/skipped_tests.tbl @@ -820,9 +820,13 @@ tests/third_party/cupy/math_tests/test_sumprod.py::TestNansumNanprodLong_param_9 tests/third_party/cupy/math_tests/test_sumprod.py::TestNansumNanprodLong_param_9_{axis=0, func='nanprod', keepdims=True, shape=(2, 3, 4), transpose_axes=False}::test_nansum_axis_transposed tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_1dim tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_1dim_with_discont +tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_1dim_with_period +tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_1dim_with_discont_and_period tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_2dim_with_axis tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_2dim_with_discont tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_2dim_without_axis +tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_2dim_with_period +tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_2dim_with_discont_and_period tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_0_{a_shape=(), b_shape=(), shape=(4, 3, 2)}::test_beta tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_1_{a_shape=(), b_shape=(), shape=(3, 2)}::test_beta tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_2_{a_shape=(), b_shape=(3, 2), shape=(4, 3, 2)}::test_beta diff --git a/tests/skipped_tests_gpu.tbl b/tests/skipped_tests_gpu.tbl index 576ea0814f9..d63f4cd7422 100644 --- a/tests/skipped_tests_gpu.tbl +++ b/tests/skipped_tests_gpu.tbl @@ -948,9 +948,13 @@ tests/third_party/cupy/math_tests/test_sumprod.py::TestNansumNanprodLong_param_9 tests/third_party/cupy/math_tests/test_sumprod.py::TestNansumNanprodLong_param_9_{axis=0, func='nanprod', keepdims=True, shape=(2, 3, 4), transpose_axes=False}::test_nansum_axis_transposed tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_1dim tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_1dim_with_discont +tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_1dim_with_period +tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_1dim_with_discont_and_period tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_2dim_with_axis tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_2dim_with_discont tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_2dim_without_axis +tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_2dim_with_period +tests/third_party/cupy/math_tests/test_trigonometric.py::TestUnwrap::test_unwrap_2dim_with_discont_and_period tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_0_{a_shape=(), b_shape=(), shape=(4, 3, 2)}::test_beta tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_1_{a_shape=(), b_shape=(), shape=(3, 2)}::test_beta tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_2_{a_shape=(), b_shape=(3, 2), shape=(4, 3, 2)}::test_beta diff --git a/tests/test_umath.py b/tests/test_umath.py index e2086e77c4b..8dd98244ccb 100644 --- a/tests/test_umath.py +++ b/tests/test_umath.py @@ -1,6 +1,10 @@ import numpy import pytest -from numpy.testing import assert_allclose, assert_array_equal +from numpy.testing import ( + assert_allclose, + assert_array_almost_equal, + assert_array_equal, +) import dpnp @@ -98,20 +102,64 @@ def test_umaths(test_cases): class TestSin: - def test_sin_ordinary(self): - array_data = numpy.arange(10) - out = numpy.empty(10, dtype=numpy.float64) + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_complex=True) + ) + def test_sin(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=dpnp.float64) - dp_out = dpnp.array(out, dtype=dpnp.float64) + dp_out_dtype = dpnp.float32 + if has_support_aspect64() and dtype != dpnp.float32: + dp_out_dtype = dpnp.float64 + + dp_out = dpnp.array(np_out, dtype=dp_out_dtype) + dp_array = dpnp.array(np_array, dtype=dtype) result = dpnp.sin(dp_array, out=dp_out) # original - np_array = numpy.array(array_data, dtype=numpy.float64) - expected = numpy.sin(np_array, out=out) + expected = numpy.sin(np_array, out=np_out) - assert_array_equal(expected, result) + precision = numpy.finfo(dtype=result.dtype).precision + assert_array_almost_equal(expected, result.asnumpy(), decimal=precision) + + @pytest.mark.parametrize("dtype", get_complex_dtypes()) + def test_sin_complex(self, dtype): + np_array = numpy.arange(10, 20, dtype=dtype) + np_out = numpy.empty(10, dtype=numpy.complex128) + + # DPNP + 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.sin(dp_array, out=dp_out) + + # original + expected = numpy.sin(np_array, out=np_out) + + precision = numpy.finfo(dtype=result.dtype).precision + assert_array_almost_equal(expected, result.asnumpy(), decimal=precision) + + @pytest.mark.usefixtures("suppress_divide_numpy_warnings") + @pytest.mark.skipif( + not has_support_aspect16(), reason="No fp16 support by device" + ) + def test_sin_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.sin(dp_array, out=dp_out) + + # original + expected = numpy.sin(np_array, out=np_out) + assert_allclose(expected, result) @pytest.mark.parametrize( "dtype", @@ -119,38 +167,82 @@ def test_sin_ordinary(self): ids=["numpy.float32", "numpy.int64", "numpy.int32"], ) def test_invalid_dtype(self, dtype): - dp_array = dpnp.arange(10, dtype=dpnp.float64) + dp_array = dpnp.arange(10, dtype=dpnp.complex64) dp_out = dpnp.empty(10, dtype=dtype) - with pytest.raises(ValueError): + with pytest.raises(TypeError): dpnp.sin(dp_array, out=dp_out) @pytest.mark.parametrize( "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) def test_invalid_shape(self, shape): - dp_array = dpnp.arange(10, dtype=dpnp.float64) - dp_out = dpnp.empty(shape, dtype=dpnp.float64) + dp_array = dpnp.arange(10) + dp_out = dpnp.empty(shape, dtype=dp_array.dtype) - with pytest.raises(ValueError): + with pytest.raises(TypeError): dpnp.sin(dp_array, out=dp_out) class TestCos: - def test_cos(self): - array_data = numpy.arange(10) - out = numpy.empty(10, dtype=numpy.float64) + @pytest.mark.parametrize( + "dtype", get_all_dtypes(no_bool=True, no_complex=True) + ) + def test_cos(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=dpnp.float64) - dp_out = dpnp.array(out, dtype=dpnp.float64) + dp_out_dtype = dpnp.float32 + if has_support_aspect64() and dtype != dpnp.float32: + dp_out_dtype = dpnp.float64 + + dp_out = dpnp.array(np_out, dtype=dp_out_dtype) + dp_array = dpnp.array(np_array, dtype=dtype) result = dpnp.cos(dp_array, out=dp_out) # original - np_array = numpy.array(array_data, dtype=numpy.float64) - expected = numpy.cos(np_array, out=out) + expected = numpy.cos(np_array, out=np_out) - assert_array_equal(expected, result) + precision = numpy.finfo(dtype=result.dtype).precision + assert_array_almost_equal(expected, result.asnumpy(), decimal=precision) + + @pytest.mark.parametrize("dtype", get_complex_dtypes()) + def test_cos_complex(self, dtype): + np_array = numpy.arange(10, 20, dtype=dtype) + np_out = numpy.empty(10, dtype=numpy.complex128) + + # DPNP + 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.cos(dp_array, out=dp_out) + + # original + expected = numpy.cos(np_array, out=np_out) + + precision = numpy.finfo(dtype=result.dtype).precision + assert_array_almost_equal(expected, result.asnumpy(), decimal=precision) + + @pytest.mark.usefixtures("suppress_divide_numpy_warnings") + @pytest.mark.skipif( + not has_support_aspect16(), reason="No fp16 support by device" + ) + def test_cos_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.cos(dp_array, out=dp_out) + + # original + expected = numpy.cos(np_array, out=np_out) + assert_allclose(expected, result) @pytest.mark.parametrize( "dtype", @@ -158,20 +250,20 @@ def test_cos(self): ids=["numpy.float32", "numpy.int64", "numpy.int32"], ) def test_invalid_dtype(self, dtype): - dp_array = dpnp.arange(10, dtype=dpnp.float64) + dp_array = dpnp.arange(10, dtype=dpnp.complex64) dp_out = dpnp.empty(10, dtype=dtype) - with pytest.raises(ValueError): + with pytest.raises(TypeError): dpnp.cos(dp_array, out=dp_out) @pytest.mark.parametrize( "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) def test_invalid_shape(self, shape): - dp_array = dpnp.arange(10, dtype=dpnp.float64) - dp_out = dpnp.empty(shape, dtype=dpnp.float64) + dp_array = dpnp.arange(10) + dp_out = dpnp.empty(shape, dtype=dp_array.dtype) - with pytest.raises(ValueError): + with pytest.raises(TypeError): dpnp.cos(dp_array, out=dp_out) @@ -194,7 +286,9 @@ def test_log(self, dtype): # original expected = numpy.log(np_array, out=np_out) - assert_allclose(expected, result) + + precision = numpy.finfo(dtype=result.dtype).precision + assert_array_almost_equal(expected, result.asnumpy(), decimal=precision) @pytest.mark.parametrize("dtype", get_complex_dtypes()) def test_log_complex(self, dtype): @@ -212,7 +306,9 @@ def test_log_complex(self, dtype): # original expected = numpy.log(np_array, out=np_out) - assert_allclose(expected, result) + + precision = numpy.finfo(dtype=result.dtype).precision + assert_array_almost_equal(expected, result.asnumpy(), decimal=precision) @pytest.mark.usefixtures("suppress_divide_numpy_warnings") @pytest.mark.skipif( diff --git a/tests/third_party/cupy/math_tests/test_trigonometric.py b/tests/third_party/cupy/math_tests/test_trigonometric.py index 41064306769..60533205d93 100644 --- a/tests/third_party/cupy/math_tests/test_trigonometric.py +++ b/tests/third_party/cupy/math_tests/test_trigonometric.py @@ -3,10 +3,9 @@ from tests.third_party.cupy import testing -@testing.gpu class TestTrigonometric(unittest.TestCase): @testing.for_all_dtypes(no_complex=True) - @testing.numpy_cupy_allclose(atol=1e-5) + @testing.numpy_cupy_allclose(atol=1e-5, type_check=False) def check_unary(self, name, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) return getattr(xp, name)(a) @@ -55,7 +54,7 @@ def test_rad2deg(self): self.check_unary("rad2deg") -@testing.gpu +@testing.with_requires("numpy>=1.21.0") class TestUnwrap(unittest.TestCase): @testing.for_all_dtypes(no_complex=True) @testing.numpy_cupy_allclose() @@ -69,6 +68,18 @@ def test_unwrap_1dim_with_discont(self, xp, dtype): a = testing.shaped_random((5,), xp, dtype) return xp.unwrap(a, discont=1.0) + @testing.for_all_dtypes(no_complex=True) + @testing.numpy_cupy_allclose() + def test_unwrap_1dim_with_period(self, xp, dtype): + a = testing.shaped_random((5,), xp, dtype) + return xp.unwrap(a, period=1.2) + + @testing.for_all_dtypes(no_complex=True) + @testing.numpy_cupy_allclose() + def test_unwrap_1dim_with_discont_and_period(self, xp, dtype): + a = testing.shaped_random((5,), xp, dtype) + return xp.unwrap(a, discont=1.0, period=1.2) + @testing.for_all_dtypes(no_complex=True) @testing.numpy_cupy_allclose() def test_unwrap_2dim_without_axis(self, xp, dtype): @@ -86,3 +97,15 @@ def test_unwrap_2dim_with_axis(self, xp, dtype): def test_unwrap_2dim_with_discont(self, xp, dtype): a = testing.shaped_random((4, 5), xp, dtype) return xp.unwrap(a, discont=5.0, axis=1) + + @testing.for_all_dtypes(no_complex=True) + @testing.numpy_cupy_allclose() + def test_unwrap_2dim_with_period(self, xp, dtype): + a = testing.shaped_random((4, 5), xp, dtype) + return xp.unwrap(a, axis=1, period=4.5) + + @testing.for_all_dtypes(no_complex=True) + @testing.numpy_cupy_allclose() + def test_unwrap_2dim_with_discont_and_period(self, xp, dtype): + a = testing.shaped_random((4, 5), xp, dtype) + return xp.unwrap(a, discont=5.0, axis=1, period=4.5) From 349f8369ea4c677003b3ca6cef3a74ba5f7459cb Mon Sep 17 00:00:00 2001 From: Anton <100830759+antonwolfy@users.noreply.github.com> Date: Fri, 7 Jul 2023 20:15:07 +0200 Subject: [PATCH 2/3] Update dpnp/backend/extensions/vm/vm_py.cpp Co-authored-by: vlad-perevezentsev --- 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 0de88cb8b41..3cd32c9a6c4 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -97,7 +97,7 @@ PYBIND11_MODULE(_vm_impl, m) }; m.def("_cos", cos_pyapi, "Call `cos` function from OneMKL VM library to compute " - "natural logarithm of vector elements", + "cosine of vector elements", py::arg("sycl_queue"), py::arg("src"), py::arg("dst"), py::arg("depends") = py::list()); From 27ed8d3c0170f1309cdde6e6ec8c6a5b68b76994 Mon Sep 17 00:00:00 2001 From: Anton <100830759+antonwolfy@users.noreply.github.com> Date: Fri, 7 Jul 2023 20:15:13 +0200 Subject: [PATCH 3/3] Update dpnp/backend/extensions/vm/vm_py.cpp Co-authored-by: vlad-perevezentsev --- 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 3cd32c9a6c4..0e26d310a2d 100644 --- a/dpnp/backend/extensions/vm/vm_py.cpp +++ b/dpnp/backend/extensions/vm/vm_py.cpp @@ -153,7 +153,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());