From 350391d0c41c74bd976ca5464b9902657befec0c Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Fri, 3 Feb 2023 05:39:40 -0600 Subject: [PATCH 1/7] dpnp.add() doesn't work properly with a scalar --- tests/skipped_tests.tbl | 5 ++--- tests/skipped_tests_gpu.tbl | 33 ++++----------------------------- 2 files changed, 6 insertions(+), 32 deletions(-) diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl index 53bdec8af0a4..c4e1a2b48d44 100644 --- a/tests/skipped_tests.tbl +++ b/tests/skipped_tests.tbl @@ -765,16 +765,15 @@ tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_para tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_547_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='remainder', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_549_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='mod', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticModf::test_modf -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_0_{name='reciprocal', nargs=1}::test_raises_with_numpy_input + tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_10_{name='remainder', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_11_{name='mod', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_1_{name='angle', nargs=1}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_4_{name='divide', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_5_{name='power', nargs=2}::test_raises_with_numpy_input -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_6_{name='subtract', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_7_{name='true_divide', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_8_{name='floor_divide', nargs=2}::test_raises_with_numpy_input -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_9_{name='fmod', nargs=2}::test_raises_with_numpy_input + tests/third_party/cupy/math_tests/test_arithmetic.py::TestBoolSubtract_param_3_{shape=(), xp=dpnp}::test_bool_subtract tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_logaddexp tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_logaddexp2 diff --git a/tests/skipped_tests_gpu.tbl b/tests/skipped_tests_gpu.tbl index af2dbd783a4e..2ae9c0fc4e3b 100644 --- a/tests/skipped_tests_gpu.tbl +++ b/tests/skipped_tests_gpu.tbl @@ -91,18 +91,7 @@ tests/third_party/cupy/indexing_tests/test_insert.py::TestDiagIndicesInvalidValu tests/third_party/cupy/indexing_tests/test_insert.py::TestDiagIndicesFrom_param_0_{shape=(3, 3)}::test_diag_indices_from tests/third_party/cupy/indexing_tests/test_insert.py::TestDiagIndicesFrom_param_1_{shape=(0, 0)}::test_diag_indices_from tests/third_party/cupy/indexing_tests/test_insert.py::TestDiagIndicesFrom_param_2_{shape=(2, 2, 2)}::test_diag_indices_from -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_295_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_303_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int64), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_375_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_383_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int64), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_439_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_447_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0., 1., 2.], [3., 4., 5.]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_455_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_463_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int64), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_519_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_527_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[0., 1., 2.], [3., 4., 5.]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_535_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_543_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int64), dtype=float64, name='floor_divide', use_dtype=False}::test_binary + tests/third_party/cupy/math_tests/test_sumprod.py::TestSumprod::test_external_prod_all tests/third_party/cupy/math_tests/test_sumprod.py::TestSumprod::test_external_prod_axis tests/third_party/cupy/math_tests/test_sumprod.py::TestSumprod::test_external_sum_all @@ -969,34 +958,20 @@ tests/third_party/cupy/manipulation_tests/test_tiling.py::TestTile_param_4_{reps tests/third_party/cupy/manipulation_tests/test_tiling.py::TestTile_param_5_{reps=(2, 3, 4, 5)}::test_array_tile tests/third_party/cupy/manipulation_tests/test_transpose.py::TestTranspose::test_moveaxis_invalid5_2 tests/third_party/cupy/manipulation_tests/test_transpose.py::TestTranspose::test_moveaxis_invalid5_3 -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_279_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_287_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0., 1., 2.], [3., 4., 5.]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_295_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_303_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_359_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_367_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[0., 1., 2.], [3., 4., 5.]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_375_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_383_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_439_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_447_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0., 1., 2.], [3., 4., 5.]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary + tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_455_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_463_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_519_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_527_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0., 1., 2.], [3., 4., 5.]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_535_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_543_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticModf::test_modf -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_0_{name='reciprocal', nargs=1}::test_raises_with_numpy_input + tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_10_{name='remainder', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_11_{name='mod', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_1_{name='angle', nargs=1}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_4_{name='divide', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_5_{name='power', nargs=2}::test_raises_with_numpy_input -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_6_{name='subtract', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_7_{name='true_divide', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_8_{name='floor_divide', nargs=2}::test_raises_with_numpy_input -tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_9_{name='fmod', nargs=2}::test_raises_with_numpy_input -tests/third_party/cupy/math_tests/test_arithmetic.py::TestBoolSubtract_param_3_{shape=(), xp=dpnp}::test_bool_subtract + tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_logaddexp tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_logaddexp2 tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_copysign_float From f768a5145914275215645c2b48ee300acf6f30a6 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Tue, 7 Feb 2023 07:24:44 -0600 Subject: [PATCH 2/7] dpnp.subtract() doesn't work properly with a scalar --- .../include/dpnp_gen_2arg_3type_tbl.hpp | 6 +- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 39 ++------ dpnp/dpnp_array.py | 4 +- dpnp/dpnp_iface_mathematical.py | 83 +++++++++-------- tests/skipped_tests.tbl | 3 +- tests/skipped_tests_gpu.tbl | 4 - tests/test_arraycreation.py | 4 +- tests/test_mathematical.py | 89 ++++++++++++++----- .../cupy/creation_tests/test_from_data.py | 1 + .../cupy/math_tests/test_arithmetic.py | 48 +++++----- 10 files changed, 148 insertions(+), 133 deletions(-) diff --git a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp index 33f5e0d19a46..6d8b0a49340e 100644 --- a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp +++ b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp @@ -184,9 +184,9 @@ MACRO_2ARG_3TYPES_OP(dpnp_power_c, MACRO_2ARG_3TYPES_OP(dpnp_subtract_c, input1_elem - input2_elem, - nullptr, - std::false_type, + sycl::sub_sat(x1, x2), + MACRO_UNPACK_TYPES(int, long), oneapi::mkl::vm::sub, - MACRO_UNPACK_TYPES(float, double)) + MACRO_UNPACK_TYPES(float, double, std::complex, std::complex)) #undef MACRO_2ARG_3TYPES_OP diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 32097d321a71..31da58c98ff0 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -1193,6 +1193,12 @@ static void func_map_elemwise_2arg_3type_core(func_map_t& fmap) func_type_map_t::find_type, func_type_map_t::find_type>}), ...); + ((fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][FT1][FTs] = + {populate_func_types(), + (void*)dpnp_subtract_c_ext()>, + func_type_map_t::find_type, + func_type_map_t::find_type>}), + ...); } template @@ -1878,39 +1884,6 @@ static void func_map_init_elemwise_2arg_3type(func_map_t& fmap) fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_DBL][eft_DBL] = { eft_DBL, (void*)dpnp_subtract_c_default}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_INT][eft_INT] = { - eft_INT, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_INT][eft_LNG] = { - eft_LNG, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_INT][eft_FLT] = { - eft_DBL, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_INT][eft_DBL] = { - eft_DBL, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_LNG][eft_INT] = { - eft_LNG, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_LNG][eft_FLT] = { - eft_DBL, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_LNG][eft_DBL] = { - eft_DBL, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_FLT][eft_INT] = { - eft_DBL, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_FLT][eft_LNG] = { - eft_DBL, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_FLT][eft_DBL] = { - eft_DBL, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_DBL][eft_INT] = { - eft_DBL, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_DBL][eft_LNG] = { - eft_DBL, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_DBL][eft_FLT] = { - eft_DBL, (void*)dpnp_subtract_c_ext}; - fmap[DPNPFuncName::DPNP_FN_SUBTRACT_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void*)dpnp_subtract_c_ext}; - func_map_elemwise_2arg_3type_helper(fmap); return; diff --git a/dpnp/dpnp_array.py b/dpnp/dpnp_array.py index a60464583cab..2843a044b044 100644 --- a/dpnp/dpnp_array.py +++ b/dpnp/dpnp_array.py @@ -270,7 +270,9 @@ def __rmul__(self, other): # '__rpow__', # '__rrshift__', # '__rshift__', - # '__rsub__', + + def __rsub__(self, other): + return dpnp.subtract(other, self) def __rtruediv__(self, other): return dpnp.true_divide(other, self) diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index e254e916b846..41657fb2d593 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -215,7 +215,7 @@ def add(x1, if x1_desc and x2_desc: return dpnp_add(x1_desc, x2_desc, dtype=dtype, out=out, where=where).get_pyobj() - return call_origin(numpy.add, x1, x2, dtype=dtype, out=out, where=where, **kwargs) + return call_origin(numpy.add, x1, x2, out=out, where=where, dtype=dtype, subok=subok, **kwargs) def around(x1, decimals=0, out=None): @@ -1145,7 +1145,7 @@ def multiply(x1, if x1_desc and x2_desc: return dpnp_multiply(x1_desc, x2_desc, dtype=dtype, out=out, where=where).get_pyobj() - return call_origin(numpy.multiply, x1, x2, dtype=dtype, out=out, where=where, **kwargs) + return call_origin(numpy.multiply, x1, x2, out=out, where=where, dtype=dtype, subok=subok, **kwargs) def nancumprod(x1, **kwargs): @@ -1520,60 +1520,67 @@ def sign(x1, **kwargs): return call_origin(numpy.sign, x1, **kwargs) -def subtract(x1, x2, dtype=None, out=None, where=True, **kwargs): +def subtract(x1, + x2, + /, + out=None, + *, + where=True, + dtype=None, + subok=True, + **kwargs): """ Subtract arguments, element-wise. For full documentation refer to :obj:`numpy.subtract`. + Returns + ------- + y : dpnp.ndarray + The difference of `x1` and `x2`, element-wise. + Limitations ----------- - Parameters ``x1`` and ``x2`` are supported as either :obj:`dpnp.ndarray` or scalar. - Parameters ``dtype``, ``out`` and ``where`` are supported with their default values. + Parameters `x1` and `x2` are supported as either :class:`dpnp.ndarray` or scalar, + but not both (at least either `x1` or `x2` should be as :class:`dpnp.ndarray`). + Parameters `out`, `where`, `dtype` and `subok` are supported with their default values. Keyword arguments ``kwargs`` are currently unsupported. - Otherwise the functions will be executed sequentially on CPU. + Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. Example ------- - >>> import dpnp as np - >>> result = np.subtract(np.array([4, 3]), np.array([2, 7])) - >>> [x for x in result] + >>> import dpnp as dp + >>> result = dp.subtract(dp.array([4, 3]), dp.array([2, 7])) + >>> print(result) [2, -4] """ - x1_is_scalar = dpnp.isscalar(x1) - x2_is_scalar = dpnp.isscalar(x2) - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_strides=False, copy_when_nondefault_queue=False) - x2_desc = dpnp.get_dpnp_descriptor(x2, copy_when_strides=False, copy_when_nondefault_queue=False) + if out is not None: + pass + elif where is not True: + pass + elif dtype is not None: + pass + elif subok is not True: + pass + elif dpnp.isscalar(x1) and dpnp.isscalar(x2): + # at least either x1 or x2 has to be an array + pass + else: + # get a common queue to copy data from the host into a device if any input is scalar + queue = get_common_allocation_queue([x1, x2]) if dpnp.isscalar(x1) or dpnp.isscalar(x2) else None - if x1_desc and x2_desc and not kwargs: - if not x1_desc and not x1_is_scalar: - pass - elif not x2_desc and not x2_is_scalar: - pass - elif x1_is_scalar and x2_is_scalar: - pass - elif x1_desc and x1_desc.ndim == 0: - pass - elif x1_desc and x1_desc.dtype == dpnp.bool: - pass - elif x2_desc and x2_desc.ndim == 0: - pass - elif x2_desc and x2_desc.dtype == dpnp.bool: - pass - elif dtype is not None: - pass - elif out is not None: - pass - elif not where: - pass - else: - out_desc = dpnp.get_dpnp_descriptor(out, copy_when_nondefault_queue=False) if out is not None else None - return dpnp_subtract(x1_desc, x2_desc, dtype=dtype, out=out_desc, where=where).get_pyobj() + x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_strides=False, copy_when_nondefault_queue=False, alloc_queue=queue) + x2_desc = dpnp.get_dpnp_descriptor(x2, copy_when_strides=False, copy_when_nondefault_queue=False, alloc_queue=queue) + if x1_desc and x2_desc: + if x1_desc.dtype == x2_desc.dtype == dpnp.bool: + raise TypeError("DPNP boolean subtract, the `-` operator, is not supported, " + "use the bitwise_xor, the `^` operator, or the logical_xor function instead.") + return dpnp_subtract(x1_desc, x2_desc, dtype=dtype, out=out, where=where).get_pyobj() - return call_origin(numpy.subtract, x1, x2, dtype=dtype, out=out, where=where, **kwargs) + return call_origin(numpy.subtract, x1, x2, out=out, where=where, dtype=dtype, subok=subok, **kwargs) def sum(x1, axis=None, dtype=None, out=None, keepdims=False, initial=None, where=True): diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl index c4e1a2b48d44..25d1fd1bc0f5 100644 --- a/tests/skipped_tests.tbl +++ b/tests/skipped_tests.tbl @@ -389,7 +389,7 @@ tests/third_party/cupy/creation_tests/test_from_data.py::TestFromData::test_asar tests/third_party/cupy/creation_tests/test_from_data.py::TestFromData::test_ascontiguousarray_on_noncontiguous_array tests/third_party/cupy/creation_tests/test_from_data.py::TestFromData::test_asfortranarray_cuda_array_zero_dim tests/third_party/cupy/creation_tests/test_from_data.py::TestFromData::test_asfortranarray_cuda_array_zero_dim_dtype -tests/third_party/cupy/creation_tests/test_from_data.py::TestFromData::test_fromfile + tests/third_party/cupy/creation_tests/test_ranges.py::TestMeshgrid_param_0_{copy=False, indexing='xy', sparse=False}::test_meshgrid0 tests/third_party/cupy/creation_tests/test_ranges.py::TestMeshgrid_param_0_{copy=False, indexing='xy', sparse=False}::test_meshgrid1 tests/third_party/cupy/creation_tests/test_ranges.py::TestMeshgrid_param_0_{copy=False, indexing='xy', sparse=False}::test_meshgrid2 @@ -774,7 +774,6 @@ tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNu tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_7_{name='true_divide', nargs=2}::test_raises_with_numpy_input tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticRaisesWithNumpyInput_param_8_{name='floor_divide', nargs=2}::test_raises_with_numpy_input -tests/third_party/cupy/math_tests/test_arithmetic.py::TestBoolSubtract_param_3_{shape=(), xp=dpnp}::test_bool_subtract tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_logaddexp tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_logaddexp2 tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_copysign_float diff --git a/tests/skipped_tests_gpu.tbl b/tests/skipped_tests_gpu.tbl index 2ae9c0fc4e3b..be5a68c77d53 100644 --- a/tests/skipped_tests_gpu.tbl +++ b/tests/skipped_tests_gpu.tbl @@ -18,7 +18,6 @@ tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-conjugate-data2] tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-copy-data3] tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-cumprod-data4] tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-cumsum-data5] -tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-diff-data6] tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-ediff1d-data7] tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-fabs-data8] tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-floor-data9] @@ -29,11 +28,9 @@ tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-conjugate-data2] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-copy-data3] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-cumprod-data4] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-cumsum-data5] -tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-diff-data6] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-ediff1d-data7] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-fabs-data8] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-floor-data9] -tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-gradient-data10] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-nancumprod-data11] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-nancumsum-data12] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-nanprod-data13] @@ -554,7 +551,6 @@ tests/third_party/cupy/creation_tests/test_from_data.py::TestFromData::test_asar tests/third_party/cupy/creation_tests/test_from_data.py::TestFromData::test_ascontiguousarray_on_noncontiguous_array tests/third_party/cupy/creation_tests/test_from_data.py::TestFromData::test_asfortranarray_cuda_array_zero_dim tests/third_party/cupy/creation_tests/test_from_data.py::TestFromData::test_asfortranarray_cuda_array_zero_dim_dtype -tests/third_party/cupy/creation_tests/test_from_data.py::TestFromData::test_fromfile tests/third_party/cupy/creation_tests/test_ranges.py::TestMeshgrid_param_0_{copy=False, indexing='xy', sparse=False}::test_meshgrid0 tests/third_party/cupy/creation_tests/test_ranges.py::TestMeshgrid_param_0_{copy=False, indexing='xy', sparse=False}::test_meshgrid1 diff --git a/tests/test_arraycreation.py b/tests/test_arraycreation.py index 833ea6109c3c..7216f934b3ef 100644 --- a/tests/test_arraycreation.py +++ b/tests/test_arraycreation.py @@ -109,7 +109,7 @@ def test_frombuffer(dtype): @pytest.mark.usefixtures("allow_fall_back_on_numpy") -@pytest.mark.parametrize("dtype", get_all_dtypes(no_float16=False)) +@pytest.mark.parametrize("dtype", get_all_dtypes()) def test_fromfile(dtype): with tempfile.TemporaryFile() as fh: fh.write(b"\x00\x01\x02\x03\x04\x05\x06\x07\x08") @@ -275,6 +275,7 @@ def test_tri_default_dtype(): '[[1, 2], [3, 4]]', '[[0, 1, 2], [3, 4, 5], [6, 7, 8]]', '[[0, 1, 2, 3, 4], [5, 6, 7, 8, 9]]']) +# TODO: add fixture 'allow_fall_back_on_numpy' and remove operator.index() def test_tril(m, k): a = numpy.array(m) ia = dpnp.array(a) @@ -295,6 +296,7 @@ def test_tril(m, k): '[[1, 2], [3, 4]]', '[[0, 1, 2], [3, 4, 5], [6, 7, 8]]', '[[0, 1, 2, 3, 4], [5, 6, 7, 8, 9]]']) +# TODO: add fixture 'allow_fall_back_on_numpy' and remove operator.index() def test_triu(m, k): a = numpy.array(m) ia = dpnp.array(a) diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index 70e0bd73dc50..6f7ee58c0380 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -66,7 +66,7 @@ def test_diff(array): @pytest.mark.parametrize("dtype1", get_all_dtypes()) @pytest.mark.parametrize("dtype2", get_all_dtypes()) @pytest.mark.parametrize("func", - ['add', 'multiply']) + ['add', 'multiply', 'subtract']) @pytest.mark.parametrize("data", [[[1, 2], [3, 4]]], ids=['[[1, 2], [3, 4]]']) @@ -77,9 +77,14 @@ def test_op_multiple_dtypes(dtype1, func, dtype2, data): np_b = numpy.array(data, dtype=dtype2) dpnp_b = dpnp.array(data, dtype=dtype2) - result = getattr(dpnp, func)(dpnp_a, dpnp_b) - expected = getattr(numpy, func)(np_a, np_b) - assert_array_equal(result, expected) + if func == 'subtract' and (dtype1 == dtype2 == dpnp.bool): + with pytest.raises(TypeError): + result = getattr(dpnp, func)(dpnp_a, dpnp_b) + expected = getattr(numpy, func)(np_a, np_b) + else: + result = getattr(dpnp, func)(dpnp_a, dpnp_b) + expected = getattr(numpy, func)(np_a, np_b) + assert_array_equal(result, expected) @pytest.mark.parametrize("rhs", [[[1, 2, 3], [4, 5, 6]], [2.0, 1.5, 1.0], 3, 0.3]) @@ -98,15 +103,20 @@ def array_or_scalar(xp, data, dtype=None): return xp.array(data, dtype=dtype) def _test_mathematical(self, name, dtype, lhs, rhs): - a = self.array_or_scalar(dpnp, lhs, dtype=dtype) - b = self.array_or_scalar(dpnp, rhs, dtype=dtype) - result = getattr(dpnp, name)(a, b) + a_dpnp = self.array_or_scalar(dpnp, lhs, dtype=dtype) + b_dpnp = self.array_or_scalar(dpnp, rhs, dtype=dtype) - a = self.array_or_scalar(numpy, lhs, dtype=dtype) - b = self.array_or_scalar(numpy, rhs, dtype=dtype) - expected = getattr(numpy, name)(a, b) + a_np = self.array_or_scalar(numpy, lhs, dtype=dtype) + b_np = self.array_or_scalar(numpy, rhs, dtype=dtype) - assert_allclose(result, expected, atol=1e-4) + if name == 'subtract' and not numpy.isscalar(rhs) and dtype == dpnp.bool: + with pytest.raises(TypeError): + result = getattr(dpnp, name)(a_dpnp, b_dpnp) + expected = getattr(numpy, name)(a_np, b_np) + else: + result = getattr(dpnp, name)(a_dpnp, b_dpnp) + expected = getattr(numpy, name)(a_np, b_np) + assert_allclose(result, expected, atol=1e-4) @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_add(self, dtype, lhs, rhs): @@ -166,8 +176,7 @@ def test_remainder(self, dtype, lhs, rhs): def test_power(self, dtype, lhs, rhs): self._test_mathematical('power', dtype, lhs, rhs) - @pytest.mark.usefixtures("allow_fall_back_on_numpy") - @pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True, no_complex=True)) + @pytest.mark.parametrize("dtype", get_all_dtypes()) def test_subtract(self, dtype, lhs, rhs): self._test_mathematical('subtract', dtype, lhs, rhs) @@ -177,7 +186,7 @@ def test_subtract(self, dtype, lhs, rhs): ids=['bool', 'int', 'float']) @pytest.mark.parametrize("data_type", get_all_dtypes()) @pytest.mark.parametrize("func", - ['add', 'multiply']) + ['add', 'multiply', 'subtract']) @pytest.mark.parametrize("val", [0, 1, 5], ids=['0', '1', '5']) @@ -197,22 +206,28 @@ def test_op_with_scalar(array, val, func, data_type, val_type): dpnp_a = dpnp.array(array, dtype=data_type) val_ = val_type(val) - result = getattr(dpnp, func)(dpnp_a, val_) - expected = getattr(numpy, func)(np_a, val_) - assert_array_equal(result, expected) + if func == 'subtract' and val_type == bool and data_type == dpnp.bool: + with pytest.raises(TypeError): + result = getattr(dpnp, func)(dpnp_a, val_) + expected = getattr(numpy, func)(np_a, val_) - result = getattr(dpnp, func)(val_, dpnp_a) - expected = getattr(numpy, func)(val_, np_a) - assert_array_equal(result, expected) + result = getattr(dpnp, func)(val_, dpnp_a) + expected = getattr(numpy, func)(val_, np_a) + else: + result = getattr(dpnp, func)(dpnp_a, val_) + expected = getattr(numpy, func)(np_a, val_) + assert_array_equal(result, expected) + + result = getattr(dpnp, func)(val_, dpnp_a) + expected = getattr(numpy, func)(val_, np_a) + assert_array_equal(result, expected) @pytest.mark.parametrize("shape", [(), (3, 2)], ids=['()', '(3, 2)']) -@pytest.mark.parametrize("dtype", - [numpy.float32, numpy.float64], - ids=['numpy.float32', 'numpy.float64']) -def test_multiply_scalar2(shape, dtype): +@pytest.mark.parametrize("dtype", get_all_dtypes()) +def test_multiply_scalar(shape, dtype): np_a = numpy.ones(shape, dtype=dtype) dpnp_a = dpnp.ones(shape, dtype=dtype) @@ -221,6 +236,32 @@ def test_multiply_scalar2(shape, dtype): assert_allclose(result, expected) +@pytest.mark.parametrize("shape", + [(), (3, 2)], + ids=['()', '(3, 2)']) +@pytest.mark.parametrize("dtype", get_all_dtypes()) +def test_add_scalar(shape, dtype): + np_a = numpy.ones(shape, dtype=dtype) + dpnp_a = dpnp.ones(shape, dtype=dtype) + + result = 0.5 + dpnp_a + 1.7 + expected = 0.5 + np_a + 1.7 + assert_allclose(result, expected) + + +@pytest.mark.parametrize("shape", + [(), (3, 2)], + ids=['()', '(3, 2)']) +@pytest.mark.parametrize("dtype", get_all_dtypes()) +def test_subtract_scalar(shape, dtype): + np_a = numpy.ones(shape, dtype=dtype) + dpnp_a = dpnp.ones(shape, dtype=dtype) + + result = 0.5 - dpnp_a - 1.7 + expected = 0.5 - np_a - 1.7 + assert_allclose(result, expected) + + @pytest.mark.usefixtures("allow_fall_back_on_numpy") @pytest.mark.parametrize("array", [[1, 2, 3, 4, 5], [1, 2, numpy.nan, 4, 5], diff --git a/tests/third_party/cupy/creation_tests/test_from_data.py b/tests/third_party/cupy/creation_tests/test_from_data.py index e07d927b1cf0..ce71ef311a56 100644 --- a/tests/third_party/cupy/creation_tests/test_from_data.py +++ b/tests/third_party/cupy/creation_tests/test_from_data.py @@ -454,6 +454,7 @@ def test_asfortranarray_cuda_array_zero_dim_dtype( a = xp.ones((), dtype=dtype_a) return xp.asfortranarray(a, dtype=dtype_b) + @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.numpy_cupy_array_equal() def test_fromfile(self, xp): with tempfile.TemporaryFile() as fh: diff --git a/tests/third_party/cupy/math_tests/test_arithmetic.py b/tests/third_party/cupy/math_tests/test_arithmetic.py index a53a8494707c..21068ece8749 100644 --- a/tests/third_party/cupy/math_tests/test_arithmetic.py +++ b/tests/third_party/cupy/math_tests/test_arithmetic.py @@ -146,31 +146,27 @@ def check_binary(self, xp): y = y.astype(numpy.complex64) # NumPy returns an output array of another type than DPNP when input ones have diffrent types. - if self.name in ('add', 'multiply') and xp is cupy: - if xp.isscalar(arg1) and xp.isscalar(arg2): - # If both are scalars, the result will be a scalar, so needs to convert into numpy-scalar. - y = numpy.asarray(y) - elif dtype1 != dtype2: - is_array_arg1 = not xp.isscalar(arg1) - is_array_arg2 = not xp.isscalar(arg2) - - is_int_float = lambda _x, _y: numpy.issubdtype(_x, numpy.integer) and numpy.issubdtype(_y, numpy.floating) - is_same_type = lambda _x, _y, _type: numpy.issubdtype(_x, _type) and numpy.issubdtype(_y, _type) - - if is_array_arg1 and is_array_arg2: - # If both inputs are arrays where one is of floating type and another - integer, - # NumPy will return an output array of always "float64" type, - # while DPNP will return the array of a wider type from the input arrays. - if is_int_float(dtype1, dtype2) or is_int_float(dtype2, dtype1): - y = y.astype(numpy.float64) - elif is_same_type(dtype1, dtype2, numpy.floating) or is_same_type(dtype1, dtype2, numpy.integer): - # If one input is an array and another - scalar, - # NumPy will return an output array of the same type as the inpupt array has, - # while DPNP will return the array of a wider type from the inputs (considering both array and scalar). - if is_array_arg1 and not is_array_arg2: - y = y.astype(dtype1) - elif is_array_arg2 and not is_array_arg1: - y = y.astype(dtype2) + if self.name in ('add', 'multiply', 'subtract') and xp is cupy and dtype1 != dtype2 and not self.use_dtype: + is_array_arg1 = not xp.isscalar(arg1) + is_array_arg2 = not xp.isscalar(arg2) + + is_int_float = lambda _x, _y: numpy.issubdtype(_x, numpy.integer) and numpy.issubdtype(_y, numpy.floating) + is_same_type = lambda _x, _y, _type: numpy.issubdtype(_x, _type) and numpy.issubdtype(_y, _type) + + if is_array_arg1 and is_array_arg2: + # If both inputs are arrays where one is of floating type and another - integer, + # NumPy will return an output array of always "float64" type, + # while DPNP will return the array of a wider type from the input arrays. + if is_int_float(dtype1, dtype2) or is_int_float(dtype2, dtype1): + y = y.astype(numpy.float64) + elif is_same_type(dtype1, dtype2, numpy.floating) or is_same_type(dtype1, dtype2, numpy.integer): + # If one input is an array and another - scalar, + # NumPy will return an output array of the same type as the inpupt array has, + # while DPNP will return the array of a wider type from the inputs (considering both array and scalar). + if is_array_arg1 and not is_array_arg2: + y = y.astype(dtype1) + elif is_array_arg2 and not is_array_arg1: + y = y.astype(dtype2) # NumPy returns different values (nan/inf) on division by zero # depending on the architecture. @@ -188,7 +184,6 @@ def check_binary(self, xp): @testing.gpu @testing.parameterize(*( testing.product({ - # TODO(unno): boolean subtract causes DeprecationWarning in numpy>=1.13 'arg1': [testing.shaped_arange((2, 3), numpy, dtype=d) for d in all_types ] + [0, 0.0, 2, 2.0], @@ -283,7 +278,6 @@ def test_modf(self, xp, dtype): 'xp': [numpy, cupy], 'shape': [(3, 2), (), (3, 0, 2)] })) -@pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.gpu class TestBoolSubtract(unittest.TestCase): From 97af87abafe81956fef9143f18b1023efdd65d1c Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Sat, 11 Feb 2023 03:22:31 -0600 Subject: [PATCH 3/7] USM type in operations with a scalar --- dpnp/dpnp_iface_mathematical.py | 10 ++++++---- tests/test_usm_type.py | 13 +++++++++++++ 2 files changed, 19 insertions(+), 4 deletions(-) diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py index 41657fb2d593..64886de23c02 100644 --- a/dpnp/dpnp_iface_mathematical.py +++ b/dpnp/dpnp_iface_mathematical.py @@ -1569,11 +1569,13 @@ def subtract(x1, # at least either x1 or x2 has to be an array pass else: - # get a common queue to copy data from the host into a device if any input is scalar - queue = get_common_allocation_queue([x1, x2]) if dpnp.isscalar(x1) or dpnp.isscalar(x2) else None + # get USM type and queue to copy scalar from the host memory into a USM allocation + usm_type, queue = get_usm_allocations([x1, x2]) if dpnp.isscalar(x1) or dpnp.isscalar(x2) else (None, None) - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_strides=False, copy_when_nondefault_queue=False, alloc_queue=queue) - x2_desc = dpnp.get_dpnp_descriptor(x2, copy_when_strides=False, copy_when_nondefault_queue=False, alloc_queue=queue) + x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_strides=False, copy_when_nondefault_queue=False, + alloc_usm_type=usm_type, alloc_queue=queue) + x2_desc = dpnp.get_dpnp_descriptor(x2, copy_when_strides=False, copy_when_nondefault_queue=False, + alloc_usm_type=usm_type, alloc_queue=queue) if x1_desc and x2_desc: if x1_desc.dtype == x2_desc.dtype == dpnp.bool: raise TypeError("DPNP boolean subtract, the `-` operator, is not supported, " diff --git a/tests/test_usm_type.py b/tests/test_usm_type.py index 15b853b3bfa2..062929bb3a0e 100644 --- a/tests/test_usm_type.py +++ b/tests/test_usm_type.py @@ -37,6 +37,19 @@ def test_coerced_usm_types_mul(usm_type_x, usm_type_y): assert z.usm_type == du.get_coerced_usm_type([usm_type_x, usm_type_y]) +@pytest.mark.parametrize("usm_type_x", list_of_usm_types, ids=list_of_usm_types) +@pytest.mark.parametrize("usm_type_y", list_of_usm_types, ids=list_of_usm_types) +def test_coerced_usm_types_subtract(usm_type_x, usm_type_y): + x = dp.arange(50, usm_type = usm_type_x) + y = dp.arange(50, usm_type = usm_type_y) + + z = 20 - x - y - 7.4 + + assert x.usm_type == usm_type_x + assert y.usm_type == usm_type_y + assert z.usm_type == du.get_coerced_usm_type([usm_type_x, usm_type_y]) + + @pytest.mark.parametrize( "func, args", [ From 53a8953486fea162ba892b74641d761b1bed792e Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Sat, 11 Feb 2023 04:58:46 -0600 Subject: [PATCH 4/7] Rollback excluded 'floor_divide' tests from skip scope --- tests/skipped_tests_gpu.tbl | 25 +++++++++++++++++++++++++ 1 file changed, 25 insertions(+) diff --git a/tests/skipped_tests_gpu.tbl b/tests/skipped_tests_gpu.tbl index be5a68c77d53..34d1795cc98d 100644 --- a/tests/skipped_tests_gpu.tbl +++ b/tests/skipped_tests_gpu.tbl @@ -89,6 +89,19 @@ tests/third_party/cupy/indexing_tests/test_insert.py::TestDiagIndicesFrom_param_ tests/third_party/cupy/indexing_tests/test_insert.py::TestDiagIndicesFrom_param_1_{shape=(0, 0)}::test_diag_indices_from tests/third_party/cupy/indexing_tests/test_insert.py::TestDiagIndicesFrom_param_2_{shape=(2, 2, 2)}::test_diag_indices_from +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_295_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_303_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int64), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_375_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_383_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int64), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_439_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_447_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0., 1., 2.], [3., 4., 5.]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_455_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_463_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int64), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_519_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_527_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[0., 1., 2.], [3., 4., 5.]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_535_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_543_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int64), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int64), dtype=float64, name='floor_divide', use_dtype=False}::test_binary + tests/third_party/cupy/math_tests/test_sumprod.py::TestSumprod::test_external_prod_all tests/third_party/cupy/math_tests/test_sumprod.py::TestSumprod::test_external_prod_axis tests/third_party/cupy/math_tests/test_sumprod.py::TestSumprod::test_external_sum_all @@ -955,8 +968,20 @@ tests/third_party/cupy/manipulation_tests/test_tiling.py::TestTile_param_5_{reps tests/third_party/cupy/manipulation_tests/test_transpose.py::TestTranspose::test_moveaxis_invalid5_2 tests/third_party/cupy/manipulation_tests/test_transpose.py::TestTranspose::test_moveaxis_invalid5_3 +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_279_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_287_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0., 1., 2.], [3., 4., 5.]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_295_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_303_{arg1=array([[1., 2., 3.], [4., 5., 6.]], dtype=float32), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_359_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_367_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[0., 1., 2.], [3., 4., 5.]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_375_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_383_{arg1=array([[1., 2., 3.], [4., 5., 6.]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_439_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_447_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0., 1., 2.], [3., 4., 5.]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_455_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_463_{arg1=array([[1, 2, 3], [4, 5, 6]], dtype=int32), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_519_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0., 1., 2.], [3., 4., 5.]], dtype=float32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary +tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_527_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0., 1., 2.], [3., 4., 5.]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_535_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]], dtype=int32), dtype=float64, name='floor_divide', use_dtype=False}::test_binary tests/third_party/cupy/math_tests/test_arithmetic.py::TestArithmeticBinary2_param_543_{arg1=array([[1, 2, 3], [4, 5, 6]]), arg2=array([[0, 1, 2], [3, 4, 5]]), dtype=float64, name='floor_divide', use_dtype=False}::test_binary From 497db34e7e6bffd59c262971a8fcb4950641672c Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Sat, 11 Feb 2023 15:51:15 -0600 Subject: [PATCH 5/7] Explicit vector operations instead of saturation functions --- dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp index 6d8b0a49340e..0ac79bd2d383 100644 --- a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp +++ b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp @@ -111,8 +111,8 @@ MACRO_2ARG_3TYPES_OP(dpnp_add_c, input1_elem + input2_elem, - sycl::add_sat(x1, x2), - MACRO_UNPACK_TYPES(int, long), + x1 + x2, + MACRO_UNPACK_TYPES(int, long, bool), oneapi::mkl::vm::add, MACRO_UNPACK_TYPES(float, double, std::complex, std::complex)) @@ -170,8 +170,8 @@ MACRO_2ARG_3TYPES_OP(dpnp_minimum_c, // requires multiplication shape1[10] with shape2[10,1] and result expected as shape[10,10] MACRO_2ARG_3TYPES_OP(dpnp_multiply_c, input1_elem* input2_elem, - nullptr, - std::false_type, + x1 * x2, + MACRO_UNPACK_TYPES(int, long, bool), oneapi::mkl::vm::mul, MACRO_UNPACK_TYPES(float, double, std::complex, std::complex)) @@ -184,8 +184,8 @@ MACRO_2ARG_3TYPES_OP(dpnp_power_c, MACRO_2ARG_3TYPES_OP(dpnp_subtract_c, input1_elem - input2_elem, - sycl::sub_sat(x1, x2), - MACRO_UNPACK_TYPES(int, long), + x1 - x2, + MACRO_UNPACK_TYPES(int, long, bool), oneapi::mkl::vm::sub, MACRO_UNPACK_TYPES(float, double, std::complex, std::complex)) From 205fa8f0f388792b703b947baa385c6a59aba6bb Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Mon, 13 Feb 2023 09:56:45 -0600 Subject: [PATCH 6/7] Use std::int32_t and std::int64_t types --- dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp index 0ac79bd2d383..9a3c69aee8e5 100644 --- a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp +++ b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp @@ -112,7 +112,7 @@ MACRO_2ARG_3TYPES_OP(dpnp_add_c, input1_elem + input2_elem, x1 + x2, - MACRO_UNPACK_TYPES(int, long, bool), + MACRO_UNPACK_TYPES(bool, std::int32_t, std::int64_t), oneapi::mkl::vm::add, MACRO_UNPACK_TYPES(float, double, std::complex, std::complex)) @@ -171,7 +171,7 @@ MACRO_2ARG_3TYPES_OP(dpnp_minimum_c, MACRO_2ARG_3TYPES_OP(dpnp_multiply_c, input1_elem* input2_elem, x1 * x2, - MACRO_UNPACK_TYPES(int, long, bool), + MACRO_UNPACK_TYPES(bool, std::int32_t, std::int64_t), oneapi::mkl::vm::mul, MACRO_UNPACK_TYPES(float, double, std::complex, std::complex)) @@ -185,7 +185,7 @@ MACRO_2ARG_3TYPES_OP(dpnp_power_c, MACRO_2ARG_3TYPES_OP(dpnp_subtract_c, input1_elem - input2_elem, x1 - x2, - MACRO_UNPACK_TYPES(int, long, bool), + MACRO_UNPACK_TYPES(bool, std::int32_t, std::int64_t), oneapi::mkl::vm::sub, MACRO_UNPACK_TYPES(float, double, std::complex, std::complex)) From 6273e5aa55c07c8aefff408bd94fc0ca15f1994f Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Thu, 16 Feb 2023 06:49:43 -0600 Subject: [PATCH 7/7] Tune tail's loop of kernel for the vector op --- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 103 ++++++++++---------- tests/test_arraycreation.py | 8 +- 2 files changed, 53 insertions(+), 58 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 31da58c98ff0..057e0805db6a 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -881,9 +881,9 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap) \ sycl::queue q = *(reinterpret_cast(q_ref)); \ \ - _DataType_input1* input1_data = static_cast<_DataType_input1 *>(const_cast(input1_in)); \ - _DataType_input2* input2_data = static_cast<_DataType_input2 *>(const_cast(input2_in)); \ - _DataType_output* result = static_cast<_DataType_output *>(result_out); \ + _DataType_input1* input1_data = static_cast<_DataType_input1*>(const_cast(input1_in)); \ + _DataType_input2* input2_data = static_cast<_DataType_input2*>(const_cast(input2_in)); \ + _DataType_output* result = static_cast<_DataType_output*>(result_out); \ \ bool use_broadcasting = !array_equal(input1_shape, input1_ndim, input2_shape, input2_ndim); \ \ @@ -896,8 +896,7 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap) shape_elem_type* input2_shape_offsets = new shape_elem_type[input2_ndim]; \ \ get_shape_offsets_inkernel(input2_shape, input2_ndim, input2_shape_offsets); \ - use_strides = \ - use_strides || !array_equal(input2_strides, input2_ndim, input2_shape_offsets, input2_ndim); \ + use_strides = use_strides || !array_equal(input2_strides, input2_ndim, input2_shape_offsets, input2_ndim); \ delete[] input2_shape_offsets; \ \ sycl::event event; \ @@ -907,19 +906,17 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap) { \ DPNPC_id<_DataType_input1>* input1_it; \ const size_t input1_it_size_in_bytes = sizeof(DPNPC_id<_DataType_input1>); \ - input1_it = reinterpret_cast*>(dpnp_memory_alloc_c(q_ref, \ - input1_it_size_in_bytes)); \ - new (input1_it) \ - DPNPC_id<_DataType_input1>(q_ref, input1_data, input1_shape, input1_strides, input1_ndim); \ + input1_it = \ + reinterpret_cast*>(dpnp_memory_alloc_c(q_ref, input1_it_size_in_bytes)); \ + new (input1_it) DPNPC_id<_DataType_input1>(q_ref, input1_data, input1_shape, input1_strides, input1_ndim); \ \ input1_it->broadcast_to_shape(result_shape, result_ndim); \ \ DPNPC_id<_DataType_input2>* input2_it; \ const size_t input2_it_size_in_bytes = sizeof(DPNPC_id<_DataType_input2>); \ - input2_it = reinterpret_cast*>(dpnp_memory_alloc_c(q_ref, \ - input2_it_size_in_bytes)); \ - new (input2_it) \ - DPNPC_id<_DataType_input2>(q_ref, input2_data, input2_shape, input2_strides, input2_ndim); \ + input2_it = \ + reinterpret_cast*>(dpnp_memory_alloc_c(q_ref, input2_it_size_in_bytes)); \ + new (input2_it) DPNPC_id<_DataType_input2>(q_ref, input2_data, input2_shape, input2_strides, input2_ndim); \ \ input2_it->broadcast_to_shape(result_shape, result_ndim); \ \ @@ -957,27 +954,26 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap) using usm_host_allocatorT = sycl::usm_allocator; \ \ size_t strides_size = 3 * result_ndim; \ - shape_elem_type *dev_strides_data = sycl::malloc_device(strides_size, q); \ + shape_elem_type* dev_strides_data = sycl::malloc_device(strides_size, q); \ \ /* create host temporary for packed strides managed by shared pointer */ \ - auto strides_host_packed = std::vector(strides_size, \ - usm_host_allocatorT(q)); \ + auto strides_host_packed = \ + std::vector(strides_size, usm_host_allocatorT(q)); \ \ /* packed vector is concatenation of result_strides, input1_strides and input2_strides */ \ std::copy(result_strides, result_strides + result_ndim, strides_host_packed.begin()); \ std::copy(input1_strides, input1_strides + result_ndim, strides_host_packed.begin() + result_ndim); \ std::copy(input2_strides, input2_strides + result_ndim, strides_host_packed.begin() + 2 * result_ndim); \ \ - auto copy_strides_ev = q.copy(strides_host_packed.data(), \ - dev_strides_data, \ - strides_host_packed.size()); \ + auto copy_strides_ev = \ + q.copy(strides_host_packed.data(), dev_strides_data, strides_host_packed.size()); \ \ auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { \ const size_t output_id = global_id[0]; /* for (size_t i = 0; i < result_size; ++i) */ \ { \ - const shape_elem_type *result_strides_data = &dev_strides_data[0]; \ - const shape_elem_type *input1_strides_data = &dev_strides_data[1]; \ - const shape_elem_type *input2_strides_data = &dev_strides_data[2]; \ + const shape_elem_type* result_strides_data = &dev_strides_data[0]; \ + const shape_elem_type* input1_strides_data = &dev_strides_data[1]; \ + const shape_elem_type* input2_strides_data = &dev_strides_data[2]; \ \ size_t input1_id = 0; \ size_t input2_id = 0; \ @@ -1013,8 +1009,10 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap) { \ event = __mkl_operation__(q, result_size, input1_data, input2_data, result); \ } \ - else if constexpr (none_of_both_types<_DataType_input1, _DataType_input2, \ - std::complex, std::complex>) \ + else if constexpr (none_of_both_types<_DataType_input1, \ + _DataType_input2, \ + std::complex, \ + std::complex>) \ { \ constexpr size_t lws = 64; \ constexpr unsigned int vec_sz = 8; \ @@ -1026,8 +1024,8 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap) auto kernel_parallel_for_func = [=](sycl::nd_item<1> nd_it) { \ auto sg = nd_it.get_sub_group(); \ const auto max_sg_size = sg.get_max_local_range()[0]; \ - const size_t start = vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + \ - sg.get_group_id()[0] * max_sg_size); \ + const size_t start = \ + vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + sg.get_group_id()[0] * max_sg_size); \ \ if (start + static_cast(vec_sz) * max_sg_size < result_size) \ { \ @@ -1036,6 +1034,7 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap) sycl::vec<_DataType_input2, vec_sz> x2 = \ sg.load(sycl::multi_ptr<_DataType_input2, global_space>(&input2_data[start])); \ sycl::vec<_DataType_output, vec_sz> res_vec; \ + \ if constexpr (both_types_are_same<_DataType_input1, _DataType_input2, __vec_types__>) \ { \ res_vec = __vec_operation__; \ @@ -1050,11 +1049,10 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap) } \ } \ sg.store(sycl::multi_ptr<_DataType_output, global_space>(&result[start]), res_vec); \ - \ } \ else \ { \ - for (size_t k = start; k < result_size; ++k) \ + for (size_t k = start + sg.get_local_id()[0]; k < result_size; k += max_sg_size) \ { \ const _DataType_output input1_elem = input1_data[k]; \ const _DataType_output input2_elem = input2_data[k]; \ @@ -1064,8 +1062,8 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap) }; \ \ auto kernel_func = [&](sycl::handler& cgh) { \ - sycl::stream out(65536, 128, cgh);\ - cgh.parallel_for>(\ + cgh.parallel_for< \ + class __name__##_sg_kernel<_DataType_output, _DataType_input1, _DataType_input2>>( \ sycl::nd_range<1>(gws_range, lws_range), kernel_parallel_for_func); \ }; \ event = q.submit(kernel_func); \ @@ -1078,7 +1076,6 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap) const _DataType_output input1_elem = input1_data[i]; \ const _DataType_output input2_elem = input2_data[i]; \ result[i] = __operation__; \ - \ }; \ auto kernel_func = [&](sycl::handler& cgh) { \ cgh.parallel_for>( \ @@ -1112,26 +1109,25 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap) { \ DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); \ DPCTLEventVectorRef dep_event_vec_ref = nullptr; \ - DPCTLSyclEventRef event_ref = __name__<_DataType_output, _DataType_input1, _DataType_input2>( \ - q_ref, \ - result_out, \ - result_size, \ - result_ndim, \ - result_shape, \ - result_strides, \ - input1_in, \ - input1_size, \ - input1_ndim, \ - input1_shape, \ - input1_strides, \ - input2_in, \ - input2_size, \ - input2_ndim, \ - input2_shape, \ - input2_strides, \ - where, \ - dep_event_vec_ref \ - ); \ + DPCTLSyclEventRef event_ref = \ + __name__<_DataType_output, _DataType_input1, _DataType_input2>(q_ref, \ + result_out, \ + result_size, \ + result_ndim, \ + result_shape, \ + result_strides, \ + input1_in, \ + input1_size, \ + input1_ndim, \ + input1_shape, \ + input1_strides, \ + input2_in, \ + input2_size, \ + input2_ndim, \ + input2_shape, \ + input2_strides, \ + where, \ + dep_event_vec_ref); \ DPCTLEvent_WaitAndThrow(event_ref); \ DPCTLEvent_Delete(event_ref); \ } \ @@ -1172,9 +1168,8 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap) const shape_elem_type*, \ const shape_elem_type*, \ const size_t*, \ - const DPCTLEventVectorRef) = __name__<_DataType_output, \ - _DataType_input1, \ - _DataType_input2>; + const DPCTLEventVectorRef) = \ + __name__<_DataType_output, _DataType_input1, _DataType_input2>; #include diff --git a/tests/test_arraycreation.py b/tests/test_arraycreation.py index 136a85f207c1..63435bca11f0 100644 --- a/tests/test_arraycreation.py +++ b/tests/test_arraycreation.py @@ -274,12 +274,12 @@ def test_tri_default_dtype(): '[[1, 2], [3, 4]]', '[[0, 1, 2], [3, 4, 5], [6, 7, 8]]', '[[0, 1, 2, 3, 4], [5, 6, 7, 8, 9]]']) -# TODO: add fixture 'allow_fall_back_on_numpy' and remove operator.index() +@pytest.mark.usefixtures("allow_fall_back_on_numpy") @pytest.mark.parametrize("dtype", get_all_dtypes(no_float16=False)) def test_tril(m, k, dtype): a = numpy.array(m, dtype=dtype) ia = dpnp.array(a) - expected = numpy.tril(a, k=operator.index(k)) + expected = numpy.tril(a, k=k) result = dpnp.tril(ia, k=k) assert_array_equal(expected, result) @@ -296,12 +296,12 @@ def test_tril(m, k, dtype): ids=['[[1, 2], [3, 4]]', '[[0, 1, 2], [3, 4, 5], [6, 7, 8]]', '[[0, 1, 2, 3, 4], [5, 6, 7, 8, 9]]']) -# TODO: add fixture 'allow_fall_back_on_numpy' and remove operator.index() +@pytest.mark.usefixtures("allow_fall_back_on_numpy") @pytest.mark.parametrize("dtype", get_all_dtypes(no_float16=False)) def test_triu(m, k, dtype): a = numpy.array(m, dtype=dtype) ia = dpnp.array(a) - expected = numpy.triu(a, k=operator.index(k)) + expected = numpy.triu(a, k=k) result = dpnp.triu(ia, k=k) assert_array_equal(expected, result)