diff --git a/dpnp/backend/extensions/ufunc/CMakeLists.txt b/dpnp/backend/extensions/ufunc/CMakeLists.txt index 5f892506b81c..d363910f74df 100644 --- a/dpnp/backend/extensions/ufunc/CMakeLists.txt +++ b/dpnp/backend/extensions/ufunc/CMakeLists.txt @@ -24,6 +24,7 @@ # ***************************************************************************** set(_elementwise_sources + ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/bitwise_count.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/common.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/degrees.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fabs.cpp diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/bitwise_count.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/bitwise_count.cpp new file mode 100644 index 000000000000..2fe60d2a5efe --- /dev/null +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/bitwise_count.cpp @@ -0,0 +1,137 @@ +//***************************************************************************** +// Copyright (c) 2025, Intel Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// - Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE +// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF +// THE POSSIBILITY OF SUCH DAMAGE. +//***************************************************************************** + +#include + +#include "dpctl4pybind11.hpp" + +#include "bitwise_count.hpp" +#include "kernels/elementwise_functions/bitwise_count.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 impl +{ +namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common; +namespace td_ns = dpctl::tensor::type_dispatch; + +/** + * @brief A factory to define pairs of supported types for which + * sycl::bitwise_count function is available. + * + * @tparam T Type of input vector `a` and of result vector `y`. + */ +template +struct OutputType +{ + using value_type = typename std::disjunction< + td_ns::TypeMapResultEntry, + td_ns::TypeMapResultEntry, + td_ns::TypeMapResultEntry, + td_ns::TypeMapResultEntry, + td_ns::TypeMapResultEntry, + td_ns::TypeMapResultEntry, + td_ns::TypeMapResultEntry, + td_ns::TypeMapResultEntry, + td_ns::DefaultResultEntry>::result_type; +}; + +using dpnp::kernels::bitwise_count::BitwiseCountFunctor; + +template +using ContigFunctor = + ew_cmn_ns::UnaryContigFunctor, + vec_sz, + n_vecs, + enable_sg_loadstore>; + +template +using StridedFunctor = + ew_cmn_ns::UnaryStridedFunctor>; + +using ew_cmn_ns::unary_contig_impl_fn_ptr_t; +using ew_cmn_ns::unary_strided_impl_fn_ptr_t; + +static unary_contig_impl_fn_ptr_t + bitwise_count_contig_dispatch_vector[td_ns::num_types]; +static int bitwise_count_output_typeid_vector[td_ns::num_types]; +static unary_strided_impl_fn_ptr_t + bitwise_count_strided_dispatch_vector[td_ns::num_types]; + +MACRO_POPULATE_DISPATCH_VECTORS(bitwise_count); +} // namespace impl + +void init_bitwise_count(py::module_ m) +{ + using arrayT = dpctl::tensor::usm_ndarray; + using event_vecT = std::vector; + { + impl::populate_bitwise_count_dispatch_vectors(); + using impl::bitwise_count_contig_dispatch_vector; + using impl::bitwise_count_output_typeid_vector; + using impl::bitwise_count_strided_dispatch_vector; + + auto bitwise_count_pyapi = [&](const arrayT &src, const arrayT &dst, + sycl::queue &exec_q, + const event_vecT &depends = {}) { + return py_int::py_unary_ufunc( + src, dst, exec_q, depends, bitwise_count_output_typeid_vector, + bitwise_count_contig_dispatch_vector, + bitwise_count_strided_dispatch_vector); + }; + m.def("_bitwise_count", bitwise_count_pyapi, "", py::arg("src"), + py::arg("dst"), py::arg("sycl_queue"), + py::arg("depends") = py::list()); + + auto bitwise_count_result_type_pyapi = [&](const py::dtype &dtype) { + return py_int::py_unary_ufunc_result_type( + dtype, bitwise_count_output_typeid_vector); + }; + m.def("_bitwise_count_result_type", bitwise_count_result_type_pyapi); + } +} +} // namespace dpnp::extensions::ufunc diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/bitwise_count.hpp b/dpnp/backend/extensions/ufunc/elementwise_functions/bitwise_count.hpp new file mode 100644 index 000000000000..53afc55f2709 --- /dev/null +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/bitwise_count.hpp @@ -0,0 +1,35 @@ +//***************************************************************************** +// Copyright (c) 2025, 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_bitwise_count(py::module_ m); +} // namespace dpnp::extensions::ufunc diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp index f9d179d5ca4e..8ff89a1b03b6 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp @@ -25,6 +25,7 @@ #include +#include "bitwise_count.hpp" #include "degrees.hpp" #include "fabs.hpp" #include "fix.hpp" @@ -52,6 +53,7 @@ namespace dpnp::extensions::ufunc */ void init_elementwise_functions(py::module_ m) { + init_bitwise_count(m); init_degrees(m); init_fabs(m); init_fix(m); diff --git a/dpnp/backend/kernels/elementwise_functions/bitwise_count.hpp b/dpnp/backend/kernels/elementwise_functions/bitwise_count.hpp new file mode 100644 index 000000000000..0d42f30ae9ec --- /dev/null +++ b/dpnp/backend/kernels/elementwise_functions/bitwise_count.hpp @@ -0,0 +1,59 @@ +//***************************************************************************** +// Copyright (c) 2025, 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 + +// dpctl tensor headers +#include "utils/type_utils.hpp" + +namespace dpnp::kernels::bitwise_count +{ +namespace tu_ns = dpctl::tensor::type_utils; + +template +struct BitwiseCountFunctor +{ + // is function constant for given argT + using is_constant = typename std::false_type; + // constant value, if constant + // constexpr resT constant_value = resT{}; + // is function defined for sycl::vec + using supports_vec = typename std::false_type; + // do both argT and resT support subgroup store/load operation + using supports_sg_loadstore = typename std::true_type; + + resT operator()(const argT &x) const + { + if constexpr (std::is_unsigned_v) { + return sycl::popcount(x); + } + else { + return sycl::popcount(sycl::abs(x)); + } + } +}; +} // namespace dpnp::kernels::bitwise_count diff --git a/dpnp/dpnp_iface_bitwise.py b/dpnp/dpnp_iface_bitwise.py index 71ca72fd52ad..d53defdb73c2 100644 --- a/dpnp/dpnp_iface_bitwise.py +++ b/dpnp/dpnp_iface_bitwise.py @@ -38,16 +38,18 @@ """ # pylint: disable=protected-access - +# pylint: disable=no-name-in-module import dpctl.tensor._tensor_elementwise_impl as ti import numpy +import dpnp.backend.extensions.ufunc._ufunc_impl as ufi from dpnp.dpnp_algo.dpnp_elementwise_common import DPNPBinaryFunc, DPNPUnaryFunc __all__ = [ "binary_repr", "bitwise_and", + "bitwise_count", "bitwise_invert", "bitwise_left_shift", "bitwise_not", @@ -215,6 +217,59 @@ def binary_repr(num, width=None): ) +_BITWISE_COUNT_DOCSTRING = """ +Computes the number of 1-bits in the absolute value of `x`. + +For full documentation refer to :obj:`numpy.bitwise_count`. + +Parameters +---------- +x : {dpnp.ndarray, usm_ndarray} + Input array, expected to have integer or boolean data type. +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 corresponding number of 1-bits in the input. Returns ``uint8`` for all + integer types. + +Limitations +----------- +Parameters `where` and `subok` are supported with their default values. +Keyword argument `kwargs` is currently unsupported. +Otherwise ``NotImplementedError`` exception will be raised. + +Examples +-------- +>>> import dpnp as np +>>> a = np.array(1023) +>>> np.bitwise_count(a) +array(10, dtype=uint8) + +>>> a = np.array([2**i - 1 for i in range(16)]) +>>> np.bitwise_count(a) +array([ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15], + dtype=uint8) + +""" + +bitwise_count = DPNPUnaryFunc( + "bitwise_count", + ufi._bitwise_count_result_type, + ufi._bitwise_count, + _BITWISE_COUNT_DOCSTRING, +) + + _BITWISE_OR_DOCSTRING = """ Computes the bitwise OR of the underlying binary representation of each element `x1_i` of the input array `x1` with the respective element `x2_i` diff --git a/dpnp/tests/test_bitwise.py b/dpnp/tests/test_bitwise.py index d41a4900c8e3..2ec2e7184701 100644 --- a/dpnp/tests/test_bitwise.py +++ b/dpnp/tests/test_bitwise.py @@ -2,9 +2,15 @@ import pytest from numpy.testing import assert_array_equal -import dpnp as inp +import dpnp -from .helper import assert_dtype_allclose, get_abs_array, get_integer_dtypes +from .helper import ( + assert_dtype_allclose, + get_abs_array, + get_integer_dtypes, + numpy_version, +) +from .third_party.cupy import testing @pytest.mark.parametrize( @@ -23,170 +29,68 @@ 3, ], ) -@pytest.mark.parametrize("dtype", [inp.bool] + get_integer_dtypes()) -class TestBitwise: +@pytest.mark.parametrize("dtype", [dpnp.bool] + get_integer_dtypes()) +class TestBitwiseBinary: @staticmethod - def array_or_scalar(xp, data, dtype=None): + def array_or_scalar(data, dtype=None): if numpy.isscalar(data): - if dtype == inp.bool: - return numpy.dtype(dtype).type(data) - return data - - if numpy.issubdtype(dtype, numpy.unsignedinteger): - data = xp.abs(xp.array(data)) - return xp.array(data, dtype=dtype) - - def _test_unary_int(self, name, data, dtype): - if numpy.isscalar(data): - pytest.skip("Input can't be scalar") - dp_a = self.array_or_scalar(inp, data, dtype=dtype) - result = getattr(inp, name)(dp_a) - - np_a = self.array_or_scalar(numpy, data, dtype=dtype) - expected = getattr(numpy, name)(np_a) + if dtype == dpnp.bool: + data = numpy.dtype(dtype).type(data) + return data, data - assert_array_equal(result, expected) - return (dp_a, np_a) + data = get_abs_array(data, dtype=dtype) + return data, dpnp.array(data) - def _test_binary_int(self, name, lhs, rhs, dtype): - if name in ("left_shift", "right_shift") and dtype == inp.bool: - pytest.skip("A shift operation isn't implemented for bool type") - elif numpy.isscalar(lhs) and numpy.isscalar(rhs): + def _test_binary(self, name, lhs, rhs, dtype): + if numpy.isscalar(lhs) and numpy.isscalar(rhs): pytest.skip("Both inputs can't be scalars") - dp_a = self.array_or_scalar(inp, lhs, dtype=dtype) - dp_b = self.array_or_scalar(inp, rhs, dtype=dtype) - result = getattr(inp, name)(dp_a, dp_b) + a, ia = self.array_or_scalar(lhs, dtype=dtype) + b, ib = self.array_or_scalar(rhs, dtype=dtype) - np_a = self.array_or_scalar(numpy, lhs, dtype=dtype) - np_b = self.array_or_scalar(numpy, rhs, dtype=dtype) - expected = getattr(numpy, name)(np_a, np_b) + result = getattr(dpnp, name)(ia, ib) + expected = getattr(numpy, name)(a, b) + assert_array_equal(result, expected) + iout = dpnp.empty_like(result) + result = getattr(dpnp, name)(ia, ib, out=iout) + assert result is iout assert_array_equal(result, expected) - return (dp_a, dp_b, np_a, np_b) + + return (ia, ib, a, b) def test_bitwise_and(self, lhs, rhs, dtype): - dp_a, dp_b, np_a, np_b = self._test_binary_int( - "bitwise_and", lhs, rhs, dtype - ) - assert_array_equal(dp_a & dp_b, np_a & np_b) - - if ( - not (inp.isscalar(dp_a) or inp.isscalar(dp_b)) - and dp_a.shape == dp_b.shape - ): - dp_a &= dp_b - np_a &= np_b - assert_array_equal(dp_a, np_a) + ia, ib, a, b = self._test_binary("bitwise_and", lhs, rhs, dtype) + assert_array_equal(ia & ib, a & b) def test_bitwise_or(self, lhs, rhs, dtype): - dp_a, dp_b, np_a, np_b = self._test_binary_int( - "bitwise_or", lhs, rhs, dtype - ) - assert_array_equal(dp_a | dp_b, np_a | np_b) - - if ( - not (inp.isscalar(dp_a) or inp.isscalar(dp_b)) - and dp_a.shape == dp_b.shape - ): - dp_a |= dp_b - np_a |= np_b - assert_array_equal(dp_a, np_a) + ia, ib, a, b = self._test_binary("bitwise_or", lhs, rhs, dtype) + assert_array_equal(ia | ib, a | b) def test_bitwise_xor(self, lhs, rhs, dtype): - dp_a, dp_b, np_a, np_b = self._test_binary_int( - "bitwise_xor", lhs, rhs, dtype - ) - assert_array_equal(dp_a ^ dp_b, np_a ^ np_b) - - if ( - not (inp.isscalar(dp_a) or inp.isscalar(dp_b)) - and dp_a.shape == dp_b.shape - ): - dp_a ^= dp_b - np_a ^= np_b - assert_array_equal(dp_a, np_a) - - def test_invert(self, lhs, rhs, dtype): - dp_a, np_a = self._test_unary_int("invert", lhs, dtype) - assert_array_equal(~dp_a, ~np_a) + ia, ib, a, b = self._test_binary("bitwise_xor", lhs, rhs, dtype) + assert_array_equal(ia ^ ib, a ^ b) def test_left_shift(self, lhs, rhs, dtype): - dp_a, dp_b, np_a, np_b = self._test_binary_int( - "left_shift", lhs, rhs, dtype - ) - assert_array_equal(dp_a << dp_b, np_a << np_b) - - if ( - not (inp.isscalar(dp_a) or inp.isscalar(dp_b)) - and dp_a.shape == dp_b.shape - ): - dp_a <<= dp_b - np_a <<= np_b - assert_array_equal(dp_a, np_a) + if numpy_version() >= "2.0.0": + _ = self._test_binary("bitwise_left_shift", lhs, rhs, dtype) + ia, ib, a, b = self._test_binary("left_shift", lhs, rhs, dtype) + assert_array_equal(ia << ib, a << b) def test_right_shift(self, lhs, rhs, dtype): - dp_a, dp_b, np_a, np_b = self._test_binary_int( - "right_shift", lhs, rhs, dtype - ) - assert_array_equal(dp_a >> dp_b, np_a >> np_b) - - if ( - not (inp.isscalar(dp_a) or inp.isscalar(dp_b)) - and dp_a.shape == dp_b.shape - ): - dp_a >>= dp_b - np_a >>= np_b - assert_array_equal(dp_a, np_a) - - def test_bitwise_aliase1(self, lhs, rhs, dtype): - if numpy.isscalar(lhs): - pytest.skip("Input can't be scalar") - dp_a = self.array_or_scalar(inp, lhs, dtype=dtype) - result1 = inp.invert(dp_a) - result2 = inp.bitwise_invert(dp_a) - assert_array_equal(result1, result2) - - result2 = inp.bitwise_not(dp_a) - assert_array_equal(result1, result2) - - def test_bitwise_aliase2(self, lhs, rhs, dtype): - if dtype == inp.bool: - pytest.skip("A shift operation isn't implemented for bool type") - elif numpy.isscalar(lhs) and numpy.isscalar(rhs): - pytest.skip("Both inputs can't be scalars") - - dp_a = self.array_or_scalar(inp, lhs, dtype=dtype) - dp_b = self.array_or_scalar(inp, rhs, dtype=dtype) - result1 = inp.left_shift(dp_a, dp_b) - result2 = inp.bitwise_left_shift(dp_a, dp_b) - assert_array_equal(result1, result2) - - result1 = inp.right_shift(dp_a, dp_b) - result2 = inp.bitwise_right_shift(dp_a, dp_b) - assert_array_equal(result1, result2) - - -@pytest.mark.parametrize("dtype", get_integer_dtypes()) -def test_invert_out(dtype): - low = 0 if numpy.issubdtype(dtype, numpy.unsignedinteger) else -5 - np_a = numpy.arange(low, 5, dtype=dtype) - dp_a = inp.array(np_a) + if numpy_version() >= "2.0.0": + _ = self._test_binary("bitwise_right_shift", lhs, rhs, dtype) + ia, ib, a, b = self._test_binary("right_shift", lhs, rhs, dtype) + assert_array_equal(ia >> ib, a >> b) - expected = numpy.invert(np_a) - dp_out = inp.empty(expected.shape, dtype=expected.dtype) - result = inp.invert(dp_a, out=dp_out) - assert result is dp_out - assert_dtype_allclose(result, expected) - -@pytest.mark.parametrize("dtype1", [inp.bool] + get_integer_dtypes()) -@pytest.mark.parametrize("dtype2", [inp.bool] + get_integer_dtypes()) +@pytest.mark.parametrize("dtype1", [dpnp.bool] + get_integer_dtypes()) +@pytest.mark.parametrize("dtype2", [dpnp.bool] + get_integer_dtypes()) class TestBitwiseInplace: def test_bitwise_and(self, dtype1, dtype2): a = get_abs_array([[-7, 6, -3, 2, -1], [0, -3, 4, 5, -6]], dtype=dtype1) b = get_abs_array([5, -2, 0, 1, 0], dtype=dtype2) - ia, ib = inp.array(a), inp.array(b) + ia, ib = dpnp.array(a), dpnp.array(b) a &= True ia &= True @@ -214,7 +118,7 @@ def test_bitwise_and(self, dtype1, dtype2): def test_bitwise_or(self, dtype1, dtype2): a = get_abs_array([[-7, 6, -3, 2, -1], [0, -3, 4, 5, -6]], dtype=dtype1) b = get_abs_array([5, -2, 0, 1, 0], dtype=dtype2) - ia, ib = inp.array(a), inp.array(b) + ia, ib = dpnp.array(a), dpnp.array(b) a |= False ia |= False @@ -242,7 +146,7 @@ def test_bitwise_or(self, dtype1, dtype2): def test_bitwise_xor(self, dtype1, dtype2): a = get_abs_array([[-7, 6, -3, 2, -1], [0, -3, 4, 5, -6]], dtype=dtype1) b = get_abs_array([5, -2, 0, 1, 0], dtype=dtype2) - ia, ib = inp.array(a), inp.array(b) + ia, ib = dpnp.array(a), dpnp.array(b) a ^= False ia ^= False @@ -250,16 +154,16 @@ def test_bitwise_xor(self, dtype1, dtype2): a = get_abs_array([[-7, 6, -3, 2, -1], [0, -3, 4, 5, -6]], dtype=dtype1) b = get_abs_array([5, -2, 0, 1, 0], dtype=dtype2) - ia, ib = inp.array(a), inp.array(b) - if numpy.issubdtype(dtype1, numpy.signedinteger) and numpy.issubdtype( - dtype2, numpy.uint64 + ia, ib = dpnp.array(a), dpnp.array(b) + if dpnp.issubdtype(dtype1, dpnp.signedinteger) and dpnp.issubdtype( + dtype2, dpnp.uint64 ): # For this special case, NumPy raises an error but dpnp works b = b.astype(numpy.int64) a ^= b ia ^= ib assert_array_equal(ia, a) - elif numpy.can_cast(dtype2, dtype1, casting="same_kind"): + elif dpnp.can_cast(dtype2, dtype1, casting="same_kind"): a ^= b ia ^= ib assert_array_equal(ia, a) @@ -277,7 +181,7 @@ class TestBitwiseShiftInplace: def test_bitwise_left_shift(self, dtype1, dtype2): a = get_abs_array([[-7, 6, -3, 2, -1], [0, -3, 4, 5, -6]], dtype=dtype1) b = get_abs_array([5, 2, 0, 1, 0], dtype=dtype2) - ia, ib = inp.array(a), inp.array(b) + ia, ib = dpnp.array(a), dpnp.array(b) a <<= True ia <<= True @@ -305,7 +209,7 @@ def test_bitwise_left_shift(self, dtype1, dtype2): def test_bitwise_right_shift(self, dtype1, dtype2): a = get_abs_array([[-7, 6, -3, 2, -1], [0, -3, 4, 5, -6]], dtype=dtype1) b = get_abs_array([5, 2, 0, 1, 0], dtype=dtype2) - ia, ib = inp.array(a), inp.array(b) + ia, ib = dpnp.array(a), dpnp.array(b) a >>= True ia >>= True @@ -329,3 +233,39 @@ def test_bitwise_right_shift(self, dtype1, dtype2): with pytest.raises(ValueError): ia >>= ib + + +@pytest.mark.parametrize( + "val", + [ + [[-7, -6, -5, -4, -3, -2, -1], [0, 1, 2, 3, 4, 5, 6]], + [-3, -2, -1, 0, 1, 2, 3], + ], +) +@pytest.mark.parametrize("dtype", [dpnp.bool] + get_integer_dtypes()) +class TestBitwiseUnary: + def _test_unary(self, name, data, dtype): + a = get_abs_array(data, dtype=dtype) + ia = dpnp.array(a) + + result = getattr(dpnp, name)(ia) + expected = getattr(numpy, name)(a) + assert_array_equal(result, expected) + + iout = dpnp.empty_like(result) + result = getattr(dpnp, name)(ia, out=iout) + assert result is iout + assert_array_equal(result, expected) + + return (ia, a) + + @testing.with_requires("numpy>=2.0") + def test_bitwise_count(self, val, dtype): + _ = self._test_unary("bitwise_count", val, dtype) + + def test_invert(self, val, dtype): + if numpy_version() >= "2.0.0": + _ = self._test_unary("bitwise_not", val, dtype) + _ = self._test_unary("bitwise_invert", val, dtype) + ia, a = self._test_unary("invert", val, dtype) + assert_array_equal(~ia, ~a) diff --git a/dpnp/tests/test_sycl_queue.py b/dpnp/tests/test_sycl_queue.py index 7b56bb6b03f4..5fe17437ab31 100644 --- a/dpnp/tests/test_sycl_queue.py +++ b/dpnp/tests/test_sycl_queue.py @@ -599,6 +599,34 @@ def test_2in_1out_diff_queue_but_equal_context(func, device): getattr(dpnp, func)(x1, x2) +@pytest.mark.parametrize("op", ["bitwise_count", "bitwise_not"]) +@pytest.mark.parametrize("device", valid_dev, ids=dev_ids) +def test_bitwise_op_1in(op, device): + x = dpnp.arange(-10, 10, device=device) + z = getattr(dpnp, op)(x) + + assert_sycl_queue_equal(x.sycl_queue, z.sycl_queue) + + +@pytest.mark.parametrize( + "op", + ["bitwise_and", "bitwise_or", "bitwise_xor", "left_shift", "right_shift"], +) +@pytest.mark.parametrize("device", valid_dev, ids=dev_ids) +def test_bitwise_op_2in(op, device): + x = dpnp.arange(25, device=device) + y = dpnp.arange(25, device=device)[::-1] + + z = getattr(dpnp, op)(x, y) + zx = getattr(dpnp, op)(x, 7) + zy = getattr(dpnp, op)(12, y) + + assert_sycl_queue_equal(z.sycl_queue, x.sycl_queue) + assert_sycl_queue_equal(z.sycl_queue, y.sycl_queue) + assert_sycl_queue_equal(zx.sycl_queue, x.sycl_queue) + assert_sycl_queue_equal(zy.sycl_queue, y.sycl_queue) + + @pytest.mark.parametrize("device", valid_dev, ids=dev_ids) @pytest.mark.parametrize( "shape1, shape2", diff --git a/dpnp/tests/test_umath.py b/dpnp/tests/test_umath.py index 8152af1f3ab8..526586564da6 100644 --- a/dpnp/tests/test_umath.py +++ b/dpnp/tests/test_umath.py @@ -73,19 +73,11 @@ def get_id(val): return val.__str__() -# implement missing umaths and to remove the list -new_umaths_numpy_20 = [ - "bitwise_count", # SAT-7323 -] - - @pytest.mark.usefixtures("allow_fall_back_on_numpy") @pytest.mark.usefixtures("suppress_divide_invalid_numpy_warnings") @pytest.mark.parametrize("test_cases", test_cases, ids=get_id) def test_umaths(test_cases): umath, args_str = test_cases - if umath in new_umaths_numpy_20: - pytest.skip("new umaths from numpy 2.0 are not supported yet") if umath in ["matmul", "matvec", "vecmat"]: sh = (4, 4) diff --git a/dpnp/tests/test_usm_type.py b/dpnp/tests/test_usm_type.py index 0fb563523359..48b21bfd023f 100644 --- a/dpnp/tests/test_usm_type.py +++ b/dpnp/tests/test_usm_type.py @@ -371,13 +371,22 @@ def test_logic_op_2in(op, usm_type_x, usm_type_y): assert z.usm_type == du.get_coerced_usm_type([usm_type_x, usm_type_y]) +@pytest.mark.parametrize("op", ["bitwise_count", "bitwise_not"]) +@pytest.mark.parametrize("usm_type", list_of_usm_types) +def test_bitwise_op_1in(op, usm_type): + x = dpnp.arange(-10, 10, usm_type=usm_type) + res = getattr(dpnp, op)(x) + + assert x.usm_type == res.usm_type == usm_type + + @pytest.mark.parametrize( "op", ["bitwise_and", "bitwise_or", "bitwise_xor", "left_shift", "right_shift"], ) @pytest.mark.parametrize("usm_type_x", list_of_usm_types) @pytest.mark.parametrize("usm_type_y", list_of_usm_types) -def test_bitwise_op(op, usm_type_x, usm_type_y): +def test_bitwise_op_2in(op, usm_type_x, usm_type_y): x = dpnp.arange(25, usm_type=usm_type_x) y = dpnp.arange(25, usm_type=usm_type_y)[::-1]