diff --git a/dpnp/backend/extensions/ufunc/CMakeLists.txt b/dpnp/backend/extensions/ufunc/CMakeLists.txt index 3580c29ef08d..e14e053f369f 100644 --- a/dpnp/backend/extensions/ufunc/CMakeLists.txt +++ b/dpnp/backend/extensions/ufunc/CMakeLists.txt @@ -32,7 +32,9 @@ 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/gcd.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/heaviside.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/lcm.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/logaddexp2.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/radians.cpp ) diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp index 178a3c5f6362..1b348b6dc0e6 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp @@ -32,7 +32,9 @@ #include "fmax.hpp" #include "fmin.hpp" #include "fmod.hpp" +#include "gcd.hpp" #include "heaviside.hpp" +#include "lcm.hpp" #include "logaddexp2.hpp" #include "radians.hpp" @@ -52,7 +54,9 @@ void init_elementwise_functions(py::module_ m) init_fmax(m); init_fmin(m); init_fmod(m); + init_gcd(m); init_heaviside(m); + init_lcm(m); init_logaddexp2(m); init_radians(m); } diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/gcd.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/gcd.cpp new file mode 100644 index 000000000000..ed804939455e --- /dev/null +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/gcd.cpp @@ -0,0 +1,186 @@ +//***************************************************************************** +// Copyright (c) 2024, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// maxification, 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 + +#include "dpctl4pybind11.hpp" + +#include "gcd.hpp" +#include "kernels/elementwise_functions/gcd.hpp" +#include "populate.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 td_ns = dpctl::tensor::type_dispatch; + +namespace impl +{ +namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common; + +template +struct OutputType +{ + using value_type = typename std::disjunction< // disjunction is C++17 + // feature, supported by DPC++ + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::DefaultResultEntry>::result_type; +}; + +using dpnp::kernels::gcd::GcdFunctor; + +template +using ContigFunctor = + ew_cmn_ns::BinaryContigFunctor, + vec_sz, + n_vecs, + enable_sg_loadstore>; + +template +using StridedFunctor = + ew_cmn_ns::BinaryStridedFunctor>; + +using ew_cmn_ns::binary_contig_impl_fn_ptr_t; +using ew_cmn_ns::binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t; +using ew_cmn_ns::binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t; +using ew_cmn_ns::binary_strided_impl_fn_ptr_t; + +static binary_contig_impl_fn_ptr_t gcd_contig_dispatch_table[td_ns::num_types] + [td_ns::num_types]; +static int gcd_output_typeid_table[td_ns::num_types][td_ns::num_types]; +static binary_strided_impl_fn_ptr_t + gcd_strided_dispatch_table[td_ns::num_types][td_ns::num_types]; + +MACRO_POPULATE_DISPATCH_TABLES(gcd); +} // namespace impl + +void init_gcd(py::module_ m) +{ + using arrayT = dpctl::tensor::usm_ndarray; + using event_vecT = std::vector; + { + impl::populate_gcd_dispatch_tables(); + using impl::gcd_contig_dispatch_table; + using impl::gcd_output_typeid_table; + using impl::gcd_strided_dispatch_table; + + auto gcd_pyapi = [&](const arrayT &src1, const arrayT &src2, + const arrayT &dst, sycl::queue &exec_q, + const event_vecT &depends = {}) { + return py_int::py_binary_ufunc( + src1, src2, dst, exec_q, depends, gcd_output_typeid_table, + gcd_contig_dispatch_table, gcd_strided_dispatch_table, + // no dedicated kernel for C-contig row with broadcasting + td_ns::NullPtrTable< + impl:: + binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{}, + td_ns::NullPtrTable< + impl:: + binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{}); + }; + m.def("_gcd", gcd_pyapi, "", py::arg("src1"), py::arg("src2"), + py::arg("dst"), py::arg("sycl_queue"), + py::arg("depends") = py::list()); + + auto gcd_result_type_pyapi = [&](const py::dtype &dtype1, + const py::dtype &dtype2) { + return py_int::py_binary_ufunc_result_type(dtype1, dtype2, + gcd_output_typeid_table); + }; + m.def("_gcd_result_type", gcd_result_type_pyapi); + } +} +} // namespace dpnp::extensions::ufunc diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/gcd.hpp b/dpnp/backend/extensions/ufunc/elementwise_functions/gcd.hpp new file mode 100644 index 000000000000..7b82b1ae0210 --- /dev/null +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/gcd.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 + +namespace py = pybind11; + +namespace dpnp::extensions::ufunc +{ +void init_gcd(py::module_ m); +} // namespace dpnp::extensions::ufunc diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.cpp index c4e4603460f5..9af1dcc9a6e9 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.cpp @@ -38,7 +38,6 @@ // dpctl tensor headers #include "kernels/elementwise_functions/common.hpp" -#include "kernels/elementwise_functions/maximum.hpp" #include "utils/type_dispatch.hpp" namespace dpnp::extensions::ufunc diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/lcm.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/lcm.cpp new file mode 100644 index 000000000000..d02e98236240 --- /dev/null +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/lcm.cpp @@ -0,0 +1,186 @@ +//***************************************************************************** +// Copyright (c) 2024, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// maxification, 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 + +#include "dpctl4pybind11.hpp" + +#include "kernels/elementwise_functions/lcm.hpp" +#include "lcm.hpp" +#include "populate.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 td_ns = dpctl::tensor::type_dispatch; + +namespace impl +{ +namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common; + +template +struct OutputType +{ + using value_type = typename std::disjunction< // disjunction is C++17 + // feature, supported by DPC++ + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::DefaultResultEntry>::result_type; +}; + +using dpnp::kernels::lcm::LcmFunctor; + +template +using ContigFunctor = + ew_cmn_ns::BinaryContigFunctor, + vec_sz, + n_vecs, + enable_sg_loadstore>; + +template +using StridedFunctor = + ew_cmn_ns::BinaryStridedFunctor>; + +using ew_cmn_ns::binary_contig_impl_fn_ptr_t; +using ew_cmn_ns::binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t; +using ew_cmn_ns::binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t; +using ew_cmn_ns::binary_strided_impl_fn_ptr_t; + +static binary_contig_impl_fn_ptr_t lcm_contig_dispatch_table[td_ns::num_types] + [td_ns::num_types]; +static int lcm_output_typeid_table[td_ns::num_types][td_ns::num_types]; +static binary_strided_impl_fn_ptr_t + lcm_strided_dispatch_table[td_ns::num_types][td_ns::num_types]; + +MACRO_POPULATE_DISPATCH_TABLES(lcm); +} // namespace impl + +void init_lcm(py::module_ m) +{ + using arrayT = dpctl::tensor::usm_ndarray; + using event_vecT = std::vector; + { + impl::populate_lcm_dispatch_tables(); + using impl::lcm_contig_dispatch_table; + using impl::lcm_output_typeid_table; + using impl::lcm_strided_dispatch_table; + + auto lcm_pyapi = [&](const arrayT &src1, const arrayT &src2, + const arrayT &dst, sycl::queue &exec_q, + const event_vecT &depends = {}) { + return py_int::py_binary_ufunc( + src1, src2, dst, exec_q, depends, lcm_output_typeid_table, + lcm_contig_dispatch_table, lcm_strided_dispatch_table, + // no dedicated kernel for C-contig row with broadcasting + td_ns::NullPtrTable< + impl:: + binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{}, + td_ns::NullPtrTable< + impl:: + binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{}); + }; + m.def("_lcm", lcm_pyapi, "", py::arg("src1"), py::arg("src2"), + py::arg("dst"), py::arg("sycl_queue"), + py::arg("depends") = py::list()); + + auto lcm_result_type_pyapi = [&](const py::dtype &dtype1, + const py::dtype &dtype2) { + return py_int::py_binary_ufunc_result_type(dtype1, dtype2, + lcm_output_typeid_table); + }; + m.def("_lcm_result_type", lcm_result_type_pyapi); + } +} +} // namespace dpnp::extensions::ufunc diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/lcm.hpp b/dpnp/backend/extensions/ufunc/elementwise_functions/lcm.hpp new file mode 100644 index 000000000000..7f49171afb7b --- /dev/null +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/lcm.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 + +namespace py = pybind11; + +namespace dpnp::extensions::ufunc +{ +void init_lcm(py::module_ m); +} // namespace dpnp::extensions::ufunc diff --git a/dpnp/backend/kernels/elementwise_functions/gcd.hpp b/dpnp/backend/kernels/elementwise_functions/gcd.hpp new file mode 100644 index 000000000000..e643e9559003 --- /dev/null +++ b/dpnp/backend/kernels/elementwise_functions/gcd.hpp @@ -0,0 +1,44 @@ +//***************************************************************************** +// 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 +#include + +namespace dpnp::kernels::gcd +{ +template +struct GcdFunctor +{ + using supports_sg_loadstore = typename std::true_type; + using supports_vec = typename std::false_type; + + resT operator()(const argT1 &in1, const argT2 &in2) const + { + return oneapi::dpl::gcd(in1, in2); + } +}; +} // namespace dpnp::kernels::gcd diff --git a/dpnp/backend/kernels/elementwise_functions/lcm.hpp b/dpnp/backend/kernels/elementwise_functions/lcm.hpp new file mode 100644 index 000000000000..f1a346addd46 --- /dev/null +++ b/dpnp/backend/kernels/elementwise_functions/lcm.hpp @@ -0,0 +1,44 @@ +//***************************************************************************** +// 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 +#include + +namespace dpnp::kernels::lcm +{ +template +struct LcmFunctor +{ + using supports_sg_loadstore = typename std::true_type; + using supports_vec = typename std::false_type; + + resT operator()(const argT1 &in1, const argT2 &in2) const + { + return oneapi::dpl::lcm(in1, in2); + } +}; +} // namespace dpnp::kernels::lcm diff --git a/dpnp/dpnp_algo/dpnp_elementwise_common.py b/dpnp/dpnp_algo/dpnp_elementwise_common.py index 2a8e2cc425fc..0793306a1b44 100644 --- a/dpnp/dpnp_algo/dpnp_elementwise_common.py +++ b/dpnp/dpnp_algo/dpnp_elementwise_common.py @@ -36,6 +36,7 @@ from dpnp.dpnp_array import dpnp_array __all__ = [ + "acceptance_fn_gcd_lcm", "acceptance_fn_negative", "acceptance_fn_positive", "acceptance_fn_sign", @@ -548,6 +549,18 @@ def __call__(self, x, decimals=0, out=None, dtype=None): return super().__call__(x, out=out, dtype=dtype) +def acceptance_fn_gcd_lcm( + arg1_dtype, arg2_dtype, buf1_dt, buf2_dt, res_dt, sycl_dev +): + # gcd/lcm are not defined for boolean data type + if arg1_dtype.char == "?" and arg2_dtype.char == "?": + raise ValueError( + "The function is not supported for inputs of data type bool" + ) + else: + return True + + def acceptance_fn_negative(arg_dtype, buf_dt, res_dt, sycl_dev): # negative is not defined for boolean data type if arg_dtype.char == "?": diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index 050f229e8a79..8e0faf5d79d3 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -63,6 +63,7 @@ DPNPReal, DPNPRound, DPNPUnaryFunc, + acceptance_fn_gcd_lcm, acceptance_fn_negative, acceptance_fn_positive, acceptance_fn_sign, @@ -99,9 +100,11 @@ "fmax", "fmin", "fmod", + "gcd", "gradient", "heaviside", "imag", + "lcm", "maximum", "minimum", "mod", @@ -1978,6 +1981,55 @@ def ediff1d(ary, to_end=None, to_begin=None): mkl_impl_fn="_fmod", ) +_GCD_DOCSTRING = """ +Returns the greatest common divisor of ``|x1|`` and ``|x2|``. + +For full documentation refer to :obj:`numpy.gcd`. + +Parameters +---------- +x1 : {dpnp.ndarray, usm_ndarray, scalar} + First input array, expected to have an integer data type. + Both inputs `x1` and `x2` can not be scalars at the same time. +x2 : {dpnp.ndarray, usm_ndarray, scalar} + Second input array, also expected to have an integer data type. + Both inputs `x1` and `x2` can not be scalars at the same time. +x : {dpnp.ndarray, usm_ndarray} + An array of floats to be rounded. +out : {None, dpnp.ndarray, usm_ndarray}, optional + Output array to populate. + Array must have the correct shape and the expected data type. + Default: ``None``. +order : {"C", "F", "A", "K"}, optional + Memory layout of the newly output array, if parameter `out` is ``None``. + Default: ``"K"``. + +Returns +------- +out : dpnp.ndarray + The greatest common divisor of the absolute value of the inputs. + +See Also +-------- +:obj:`dpnp.lcm` : The lowest common multiple. + +Examples +-------- +>>> import dpnp as np +>>> np.gcd(np.array(12), 20) +array(4) +>>> np.gcd(np.arange(6), 20) +array([20, 1, 2, 1, 4, 5]) +""" + +gcd = DPNPBinaryFunc( + "gcd", + ufi._gcd_result_type, + ufi._gcd, + _GCD_DOCSTRING, + acceptance_fn=acceptance_fn_gcd_lcm, +) + def gradient(f, *varargs, axis=None, edge_order=1): """ @@ -2295,6 +2347,54 @@ def gradient(f, *varargs, axis=None, edge_order=1): ) +_LCM_DOCSTRING = """ +Returns the lowest common multiple of ``|x1|`` and ``|x2|``. + +For full documentation refer to :obj:`numpy.lcm`. + +Parameters +---------- +x1 : {dpnp.ndarray, usm_ndarray, scalar} + First input array, expected to have an integer data type. + Both inputs `x1` and `x2` can not be scalars at the same time. +x2 : {dpnp.ndarray, usm_ndarray, scalar} + Second input array, also expected to have an integer data type. + Both inputs `x1` and `x2` can not be scalars at the same time. +out : {None, dpnp.ndarray, usm_ndarray}, optional + Output array to populate. + Array must have the correct shape and the expected data type. + Default: ``None``. +order : {"C", "F", "A", "K"}, optional + Memory layout of the newly output array, if parameter `out` is ``None``. + Default: ``"K"``. + +Returns +------- +out : dpnp.ndarray + The lowest common multiple of the absolute value of the inputs. + +See Also +-------- +:obj:`dpnp.gcd` : The greatest common divisor. + +Examples +-------- +>>> import dpnp as np +>>> np.lcm(np.array(12), 20) +array(60) +>>> np.lcm(np.arange(6), 20) +array([ 0, 20, 20, 60, 20, 20]) +""" + +lcm = DPNPBinaryFunc( + "lcm", + ufi._lcm_result_type, + ufi._lcm, + _LCM_DOCSTRING, + acceptance_fn=acceptance_fn_gcd_lcm, +) + + _MAXIMUM_DOCSTRING = """ Compares two input arrays `x1` and `x2` and returns a new array containing the element-wise maxima. diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index dbaab556e0c2..52677b0403bb 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -1724,6 +1724,64 @@ def test_prod_Error(self): dpnp.prod(ia, initial=6) +class TestRationalFunctions: + @pytest.mark.parametrize("func", ["gcd", "lcm"]) + @pytest.mark.parametrize("dt1", get_integer_dtypes()) + @pytest.mark.parametrize("dt2", get_integer_dtypes()) + def test_basic(self, func, dt1, dt2): + a = numpy.array([12, 120], dtype=dt1) + b = numpy.array([20, 120], dtype=dt2) + ia, ib = dpnp.array(a), dpnp.array(b) + + expected = getattr(numpy, func)(a, b) + result = getattr(dpnp, func)(ia, ib) + assert_array_equal(result, expected) + + @pytest.mark.parametrize("func", ["gcd", "lcm"]) + @pytest.mark.parametrize("dt", get_integer_dtypes()) + def test_broadcasting(self, func, dt): + a = numpy.arange(6, dtype=dt) + ia = dpnp.array(a) + b = 20 + + expected = getattr(numpy, func)(a, b) + result = getattr(dpnp, func)(ia, b) + assert_array_equal(result, expected) + + @pytest.mark.parametrize("dt", [numpy.int32, numpy.int64]) + def test_gcd_overflow(self, dt): + a = dt(numpy.iinfo(dt).min) # negative power of two + ia = dpnp.array(a) + q = -(a // 4) + + # verify that we don't overflow when taking abs(x) + # not relevant for lcm, where the result is unrepresentable anyway + expected = numpy.gcd(a, q) + result = dpnp.gcd(ia, q) + assert_array_equal(result, expected) + + def test_lcm_overflow(self): + big = numpy.int32(numpy.iinfo(numpy.int32).max // 11) + a, b = 2 * big, 5 * big + ia, ib = dpnp.array(a), dpnp.array(b) + + # verify that we don't overflow when a*b does overflow + expected = numpy.lcm(a, b) + result = dpnp.lcm(ia, ib) + assert_array_equal(result, expected) + + @pytest.mark.parametrize("func", ["gcd", "lcm"]) + @pytest.mark.parametrize("xp", [dpnp, numpy]) + def test_inf_and_nan(self, func, xp): + inf = xp.array([xp.inf]) + assert_raises((TypeError, ValueError), getattr(xp, func), inf, 1) + assert_raises((TypeError, ValueError), getattr(xp, func), 1, inf) + assert_raises((TypeError, ValueError), getattr(xp, func), xp.nan, inf) + assert_raises( + (TypeError, ValueError), getattr(xp, func), 4, float(xp.inf) + ) + + class TestRealIfClose: @pytest.mark.parametrize("dt", get_all_dtypes(no_none=True)) def test_basic(self, dt): diff --git a/tests/test_sycl_queue.py b/tests/test_sycl_queue.py index e1ae1d8e65dc..aa04497f453d 100644 --- a/tests/test_sycl_queue.py +++ b/tests/test_sycl_queue.py @@ -715,6 +715,11 @@ def test_reduce_hypot(device): [-3.0, -2.0, -1.0, 1.0, 2.0, 3.0], [2.0, 2.0, 2.0, 2.0, 2.0, 2.0], ), + pytest.param( + "gcd", + [0, 1, 2, 3, 4, 5], + [20, 20, 20, 20, 20, 20], + ), pytest.param( "gradient", [1.0, 2.0, 4.0, 7.0, 11.0, 16.0], @@ -731,6 +736,11 @@ def test_reduce_hypot(device): ), pytest.param("inner", [1.0, 2.0, 3.0], [4.0, 5.0, 6.0]), pytest.param("kron", [3.0, 4.0, 5.0], [1.0, 2.0]), + pytest.param( + "lcm", + [0, 1, 2, 3, 4, 5], + [20, 20, 20, 20, 20, 20], + ), pytest.param("logaddexp", [[-1, 2, 5, 9]], [[4, -3, 2, -8]]), pytest.param("logaddexp2", [[-1, 2, 5, 9]], [[4, -3, 2, -8]]), pytest.param( diff --git a/tests/test_usm_type.py b/tests/test_usm_type.py index 592340d6c0db..cfa65b9b1cef 100644 --- a/tests/test_usm_type.py +++ b/tests/test_usm_type.py @@ -685,6 +685,11 @@ def test_1in_1out(func, data, usm_type): pytest.param("fmax", [0.0, 1.0, 2.0], [3.0, 4.0, 5.0]), pytest.param("fmin", [0.0, 1.0, 2.0], [3.0, 4.0, 5.0]), pytest.param("fmod", [5, 3], [2, 2.0]), + pytest.param( + "gcd", + [0, 1, 2, 3, 4, 5], + [20, 20, 20, 20, 20, 20], + ), pytest.param( "gradient", [1, 2, 4, 7, 11, 16], [0.0, 1.0, 1.5, 3.5, 4.0, 6.0] ), @@ -694,6 +699,11 @@ def test_1in_1out(func, data, usm_type): ), pytest.param("inner", [1.0, 2.0, 3.0], [4.0, 5.0, 6.0]), pytest.param("kron", [3.0, 4.0, 5.0], [1.0, 2.0]), + pytest.param( + "lcm", + [0, 1, 2, 3, 4, 5], + [20, 20, 20, 20, 20, 20], + ), pytest.param("logaddexp", [[-1, 2, 5, 9]], [[4, -3, 2, -8]]), pytest.param("logaddexp2", [[-1, 2, 5, 9]], [[4, -3, 2, -8]]), pytest.param("maximum", [0.0, 1.0, 2.0], [3.0, 4.0, 5.0]), diff --git a/tests/third_party/cupy/math_tests/test_rational.py b/tests/third_party/cupy/math_tests/test_rational.py new file mode 100644 index 000000000000..6dd2a863a4ca --- /dev/null +++ b/tests/third_party/cupy/math_tests/test_rational.py @@ -0,0 +1,36 @@ +import unittest + +import pytest + +import dpnp as cupy +from tests.third_party.cupy import testing + + +class TestRational(unittest.TestCase): + @testing.for_dtypes(["?", "e", "f", "d", "F", "D"]) + def test_gcd_dtype_check(self, dtype): + a = cupy.random.randint(-10, 10, size=(10, 10)).astype(dtype) + b = cupy.random.randint(-10, 10, size=(10, 10)).astype(dtype) + with pytest.raises(ValueError): + cupy.gcd(a, b) + + @testing.for_int_dtypes(no_bool=True) + @testing.numpy_cupy_array_equal() + def test_gcd_check_boundary_cases(self, xp, dtype): + a = xp.array([0, -10, -5, 10, 410, 1, 6, 33]) + b = xp.array([0, 5, -10, -5, 20, 51, 6, 42]) + return xp.gcd(a, b) + + @testing.for_dtypes(["?", "e", "f", "d", "F", "D"]) + def test_lcm_dtype_check(self, dtype): + a = cupy.random.randint(-10, 10, size=(10, 10)).astype(dtype) + b = cupy.random.randint(-10, 10, size=(10, 10)).astype(dtype) + with pytest.raises(ValueError): + cupy.lcm(a, b) + + @testing.for_int_dtypes(no_bool=True) + @testing.numpy_cupy_array_equal() + def test_lcm_check_boundary_cases(self, xp, dtype): + a = xp.array([0, -10, -5, 10, 410, 1, 6, 33]) + b = xp.array([0, 5, -10, -5, 20, 51, 6, 42]) + return xp.lcm(a, b)