From 85dedc668f370726e0f435fe9525ab09f141c0dd Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 21 Aug 2024 17:27:42 +0200 Subject: [PATCH 1/4] Add _heaviside to _ufunc_impl extension --- dpnp/backend/extensions/ufunc/CMakeLists.txt | 1 + .../ufunc/elementwise_functions/common.cpp | 2 + .../ufunc/elementwise_functions/heaviside.cpp | 145 ++++++++++++++++++ .../ufunc/elementwise_functions/heaviside.hpp | 35 +++++ .../elementwise_functions/heaviside.hpp | 57 +++++++ dpnp/dpnp_iface_mathematical.py | 57 +++++++ 6 files changed, 297 insertions(+) create mode 100644 dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.cpp create mode 100644 dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.hpp create mode 100644 dpnp/backend/kernels/elementwise_functions/heaviside.hpp diff --git a/dpnp/backend/extensions/ufunc/CMakeLists.txt b/dpnp/backend/extensions/ufunc/CMakeLists.txt index 03e8e3eae48..3580c29ef08 100644 --- a/dpnp/backend/extensions/ufunc/CMakeLists.txt +++ b/dpnp/backend/extensions/ufunc/CMakeLists.txt @@ -32,6 +32,7 @@ set(_elementwise_sources ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmax.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmin.cpp ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmod.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/heaviside.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 6ded010d27b..178a3c5f636 100644 --- a/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp @@ -32,6 +32,7 @@ #include "fmax.hpp" #include "fmin.hpp" #include "fmod.hpp" +#include "heaviside.hpp" #include "logaddexp2.hpp" #include "radians.hpp" @@ -51,6 +52,7 @@ void init_elementwise_functions(py::module_ m) init_fmax(m); init_fmin(m); init_fmod(m); + init_heaviside(m); init_logaddexp2(m); init_radians(m); } diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.cpp new file mode 100644 index 00000000000..c4e4603460f --- /dev/null +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.cpp @@ -0,0 +1,145 @@ +//***************************************************************************** +// 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 "heaviside.hpp" +#include "kernels/elementwise_functions/heaviside.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 "kernels/elementwise_functions/maximum.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::DefaultResultEntry>::result_type; +}; + +using dpnp::kernels::heaviside::HeavisideFunctor; + +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 + heaviside_contig_dispatch_table[td_ns::num_types][td_ns::num_types]; +static int heaviside_output_typeid_table[td_ns::num_types][td_ns::num_types]; +static binary_strided_impl_fn_ptr_t + heaviside_strided_dispatch_table[td_ns::num_types][td_ns::num_types]; + +MACRO_POPULATE_DISPATCH_TABLES(heaviside); +} // namespace impl + +void init_heaviside(py::module_ m) +{ + using arrayT = dpctl::tensor::usm_ndarray; + using event_vecT = std::vector; + { + impl::populate_heaviside_dispatch_tables(); + using impl::heaviside_contig_dispatch_table; + using impl::heaviside_output_typeid_table; + using impl::heaviside_strided_dispatch_table; + + auto heaviside_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, heaviside_output_typeid_table, + heaviside_contig_dispatch_table, + heaviside_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("_heaviside", heaviside_pyapi, "", py::arg("src1"), + py::arg("src2"), py::arg("dst"), py::arg("sycl_queue"), + py::arg("depends") = py::list()); + + auto heaviside_result_type_pyapi = [&](const py::dtype &dtype1, + const py::dtype &dtype2) { + return py_int::py_binary_ufunc_result_type( + dtype1, dtype2, heaviside_output_typeid_table); + }; + m.def("_heaviside_result_type", heaviside_result_type_pyapi); + } +} +} // namespace dpnp::extensions::ufunc diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.hpp b/dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.hpp new file mode 100644 index 00000000000..ca2c6062e2c --- /dev/null +++ b/dpnp/backend/extensions/ufunc/elementwise_functions/heaviside.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_heaviside(py::module_ m); +} // namespace dpnp::extensions::ufunc diff --git a/dpnp/backend/kernels/elementwise_functions/heaviside.hpp b/dpnp/backend/kernels/elementwise_functions/heaviside.hpp new file mode 100644 index 00000000000..0700049b2ba --- /dev/null +++ b/dpnp/backend/kernels/elementwise_functions/heaviside.hpp @@ -0,0 +1,57 @@ +//***************************************************************************** +// 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 + +// dpctl tensor headers +#include "utils/math_utils.hpp" +#include "utils/type_utils.hpp" + +namespace dpnp::kernels::heaviside +{ +namespace mu_ns = dpctl::tensor::math_utils; +namespace tu_ns = dpctl::tensor::type_utils; + +template +struct HeavisideFunctor +{ + using supports_sg_loadstore = std::negation< + std::disjunction, tu_ns::is_complex>>; + using supports_vec = typename std::false_type; + + resT operator()(const argT1 &in1, const argT2 &in2) const + { + if (std::isnan(in1)) { + return in1; + } + else if (in1 == 0) { + return in2; + } + return resT(in1 > 0); + } +}; +} // namespace dpnp::kernels::heaviside diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index f3de72df3aa..a7a7508eb43 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -100,6 +100,7 @@ "fmin", "fmod", "gradient", + "heaviside", "imag", "maximum", "minimum", @@ -2179,6 +2180,62 @@ def gradient(f, *varargs, axis=None, edge_order=1): return tuple(outvals) +_HEAVISIDE_DOCSTRING = """ +Compute the Heaviside step function. + +The Heaviside step function is defined as:: + + 0 if x1 < 0 + heaviside(x1, x2) = x2 if x1 == 0 + 1 if x1 > 0 + +where `x2` is often taken to be 0.5, but 0 and 1 are also sometimes used. + +Parameters +---------- +x1 : {dpnp.ndarray, usm_ndarray, scalar} + Input values. + Both inputs `x1` and `x2` can not be scalars at the same time. +x2 : {dpnp.ndarray, usm_ndarray, scalar} + The value of the function when `x1` is ``0``. + 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 output array, element-wise Heaviside step function of `x1`. + +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([-1.5, 0, 2.0]) +>>> np.heaviside(a, 0.5) +array([0. , 0.5, 1. ]) +>>> np.heaviside(a, 1) +array([0., 1., 1.]) +""" + +heaviside = DPNPBinaryFunc( + "heaviside", + ufi._heaviside_result_type, + ufi._heaviside, + _HEAVISIDE_DOCSTRING, +) + + _IMAG_DOCSTRING = """ Computes imaginary part of each element `x_i` for input array `x`. From ca1792ce714e17de6002300e32d2cb26c2a2c97e Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 21 Aug 2024 17:32:24 +0200 Subject: [PATCH 2/4] Enable third party tests --- tests/skipped_tests.tbl | 4 ---- tests/skipped_tests_gpu.tbl | 4 ---- tests/third_party/cupy/math_tests/test_misc.py | 4 ++-- 3 files changed, 2 insertions(+), 10 deletions(-) diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl index 43ff0248058..02700e87c3f 100644 --- a/tests/skipped_tests.tbl +++ b/tests/skipped_tests.tbl @@ -19,8 +19,6 @@ tests/test_umath.py::test_umaths[('frexp', 'f')] tests/test_umath.py::test_umaths[('frexp', 'd')] tests/test_umath.py::test_umaths[('gcd', 'ii')] tests/test_umath.py::test_umaths[('gcd', 'll')] -tests/test_umath.py::test_umaths[('heaviside', 'ff')] -tests/test_umath.py::test_umaths[('heaviside', 'dd')] tests/test_umath.py::test_umaths[('lcm', 'ii')] tests/test_umath.py::test_umaths[('lcm', 'll')] tests/test_umath.py::test_umaths[('ldexp', 'fi')] @@ -193,8 +191,6 @@ tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_interp_inf_fx tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_interp_inf_x tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_interp_size1 tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_interp_inf_to_nan -tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_heaviside -tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_heaviside_nan_inf 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 diff --git a/tests/skipped_tests_gpu.tbl b/tests/skipped_tests_gpu.tbl index 32380869b43..455c2bc58a3 100644 --- a/tests/skipped_tests_gpu.tbl +++ b/tests/skipped_tests_gpu.tbl @@ -26,8 +26,6 @@ tests/test_umath.py::test_umaths[('frexp', 'f')] tests/test_umath.py::test_umaths[('frexp', 'd')] tests/test_umath.py::test_umaths[('gcd', 'ii')] tests/test_umath.py::test_umaths[('gcd', 'll')] -tests/test_umath.py::test_umaths[('heaviside', 'ff')] -tests/test_umath.py::test_umaths[('heaviside', 'dd')] tests/test_umath.py::test_umaths[('lcm', 'ii')] tests/test_umath.py::test_umaths[('lcm', 'll')] tests/test_umath.py::test_umaths[('ldexp', 'fi')] @@ -247,8 +245,6 @@ tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_interp_inf_fx tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_interp_inf_x tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_interp_size1 tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_interp_inf_to_nan -tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_heaviside -tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_heaviside_nan_inf 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 diff --git a/tests/third_party/cupy/math_tests/test_misc.py b/tests/third_party/cupy/math_tests/test_misc.py index e4eed74cd71..086f7da1728 100644 --- a/tests/third_party/cupy/math_tests/test_misc.py +++ b/tests/third_party/cupy/math_tests/test_misc.py @@ -487,7 +487,7 @@ def test_interp_inf_to_nan(self, xp, dtype_y, dtype_x): @testing.for_all_dtypes(name="dtype_2", no_bool=True, no_complex=True) @testing.for_all_dtypes(name="dtype_1", no_bool=True, no_complex=True) - @testing.numpy_cupy_array_equal() + @testing.numpy_cupy_array_equal(type_check=has_support_aspect64()) def test_heaviside(self, xp, dtype_1, dtype_2): x = testing.shaped_random((10,), xp, dtype_1) h = xp.asarray([10], dtype=dtype_2) @@ -495,7 +495,7 @@ def test_heaviside(self, xp, dtype_1, dtype_2): @testing.for_all_dtypes(name="dtype_2", no_bool=True, no_complex=True) @testing.for_float_dtypes(name="dtype_1") - @testing.numpy_cupy_array_equal() + @testing.numpy_cupy_array_equal(type_check=has_support_aspect64()) def test_heaviside_nan_inf(self, xp, dtype_1, dtype_2): x = xp.asarray([-2.0, 0.0, 3.0, xp.nan, xp.inf, -xp.inf], dtype=dtype_1) h = xp.asarray([10], dtype=dtype_2) From ef69b1885563cc2da90515e5775f779e429b525f Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 21 Aug 2024 17:33:53 +0200 Subject: [PATCH 3/4] Added CFD tests --- tests/test_sycl_queue.py | 1 + tests/test_usm_type.py | 1 + 2 files changed, 2 insertions(+) diff --git a/tests/test_sycl_queue.py b/tests/test_sycl_queue.py index 99819bfc867..cd412193eab 100644 --- a/tests/test_sycl_queue.py +++ b/tests/test_sycl_queue.py @@ -719,6 +719,7 @@ def test_reduce_hypot(device): [1.0, 2.0, 4.0, 7.0, 11.0, 16.0], [0.0, 1.0, 1.5, 3.5, 4.0, 6.0], ), + pytest.param("heaviside", [-1.5, 0, 2.0], [0.5]), pytest.param( "histogram_bin_edges", [0, 0, 0, 1, 2, 3, 3, 4, 5], diff --git a/tests/test_usm_type.py b/tests/test_usm_type.py index af81e169a3b..74f05a5bf71 100644 --- a/tests/test_usm_type.py +++ b/tests/test_usm_type.py @@ -686,6 +686,7 @@ def test_1in_1out(func, data, usm_type): pytest.param( "gradient", [1, 2, 4, 7, 11, 16], [0.0, 1.0, 1.5, 3.5, 4.0, 6.0] ), + pytest.param("heaviside", [-1.5, 0, 2.0], [1]), pytest.param( "hypot", [[1.0, 2.0, 3.0, 4.0]], [[-1.0, -2.0, -4.0, -5.0]] ), From 45c04c424a74531187f3d197a02d435bb3f7463d Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 21 Aug 2024 18:04:39 +0200 Subject: [PATCH 4/4] Added more tests to cover different use cases --- tests/test_mathematical.py | 36 ++++++++++++++++++++++++++++++++++++ 1 file changed, 36 insertions(+) diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index 0a0590a66da..7b7bd901c88 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -1230,6 +1230,42 @@ def test_return_type(self): assert type(res) is tuple +class TestHeavside: + @pytest.mark.parametrize("val", [0.5, 1.0]) + @pytest.mark.parametrize("dt", get_float_dtypes()) + def test_basic(self, val, dt): + a = numpy.array( + [[-30.0, -0.1, 0.0, 0.2], [7.5, numpy.nan, numpy.inf, -numpy.inf]], + dtype=dt, + ) + ia = dpnp.array(a) + + result = dpnp.heaviside(ia, val) + expected = numpy.heaviside(a, val) + assert_array_equal(result, expected) + + @pytest.mark.parametrize( + "a_dt", get_all_dtypes(no_none=True, no_complex=True) + ) + @pytest.mark.parametrize( + "b_dt", get_all_dtypes(no_none=True, no_complex=True) + ) + def test_both_input_as_arrays(self, a_dt, b_dt): + a = numpy.array([-1.5, 0, 2.0], dtype=a_dt) + b = numpy.array([-0, 0.5, 1.0], dtype=b_dt) + ia, ib = dpnp.array(a), dpnp.array(b) + + result = dpnp.heaviside(ia, ib) + expected = numpy.heaviside(a, b) + assert_array_equal(result, expected) + + @pytest.mark.parametrize("xp", [dpnp, numpy]) + @pytest.mark.parametrize("dt", get_complex_dtypes()) + def test_complex_dtype(self, xp, dt): + a = xp.array([-1.5, 0, 2.0], dtype=dt) + assert_raises((TypeError, ValueError), xp.heaviside, a, 0.5) + + @pytest.mark.parametrize("dtype1", get_all_dtypes()) @pytest.mark.parametrize("dtype2", get_all_dtypes()) @pytest.mark.parametrize(