From 7fad1672ce6b6a2d060929f6b1a267bb9d6958a1 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 27 Jul 2023 01:07:24 +0200 Subject: [PATCH 01/20] Update dpnp.arctan2 --- .../include/dpnp_gen_2arg_3type_tbl.hpp | 6 +- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 80 +++++++++++++++++-- 2 files changed, 75 insertions(+), 11 deletions(-) diff --git a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp index 62346c2243d..ad9a75458b7 100644 --- a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp +++ b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp @@ -114,9 +114,9 @@ MACRO_2ARG_3TYPES_OP(dpnp_add_c, std::complex)) MACRO_2ARG_3TYPES_OP(dpnp_arctan2_c, - sycl::atan2((double)input1_elem, (double)input2_elem), - nullptr, - std::false_type, + sycl::atan2(input1_elem, input2_elem), + sycl::atan2(x1, x2), + MACRO_UNPACK_TYPES(float, double), oneapi::mkl::vm::atan2, MACRO_UNPACK_TYPES(float, double)) diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 404657965ff..4caa95982b2 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -1719,25 +1719,89 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) eft_DBL, (void *)dpnp_arctan2_c_default}; fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type, int32_t, + int32_t>, + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int32_t, int32_t>}; fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_INT][eft_LNG] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type, int32_t, + int64_t>, + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int32_t, int64_t>}; fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type, int32_t, + float>, + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int32_t, float>}; fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_INT][eft_DBL] = { eft_DBL, (void *)dpnp_arctan2_c_ext}; fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_LNG][eft_INT] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type, int64_t, + int32_t>, + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int64_t, int32_t>}; fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type, int64_t, + int64_t>, + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int64_t, int64_t>}; fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type, int64_t, + float>, + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int64_t, float>}; fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_LNG][eft_DBL] = { eft_DBL, (void *)dpnp_arctan2_c_ext}; fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type, float, + int32_t>, + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + float, int32_t>}; fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type, float, + int64_t>, + get_default_floating_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + float, int64_t>}; fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_FLT][eft_FLT] = { eft_FLT, (void *)dpnp_arctan2_c_ext}; fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_FLT][eft_DBL] = { From a7d9269d24adfdc1dcdb9cc3d80f3211c76df5e5 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 27 Jul 2023 01:14:08 +0200 Subject: [PATCH 02/20] Update dpnp.copysign --- .../include/dpnp_gen_2arg_3type_tbl.hpp | 6 +- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 80 +++++++++++++++++-- 2 files changed, 75 insertions(+), 11 deletions(-) diff --git a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp index ad9a75458b7..fac16393156 100644 --- a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp +++ b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp @@ -121,9 +121,9 @@ MACRO_2ARG_3TYPES_OP(dpnp_arctan2_c, MACRO_UNPACK_TYPES(float, double)) MACRO_2ARG_3TYPES_OP(dpnp_copysign_c, - sycl::copysign((double)input1_elem, (double)input2_elem), - nullptr, - std::false_type, + sycl::copysign(input1_elem, input2_elem), + sycl::copysign(x1, x2), + MACRO_UNPACK_TYPES(float, double), oneapi::mkl::vm::copysign, MACRO_UNPACK_TYPES(float, double)) diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 4caa95982b2..73e861ef4d5 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -1849,25 +1849,89 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) eft_DBL, (void *)dpnp_copysign_c_default}; fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type, int32_t, + int32_t>, + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int32_t, int32_t>}; fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_INT][eft_LNG] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type, int32_t, + int64_t>, + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int32_t, int64_t>}; fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type, int32_t, + float>, + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int32_t, float>}; fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_INT][eft_DBL] = { eft_DBL, (void *)dpnp_copysign_c_ext}; fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_LNG][eft_INT] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type, int64_t, + int32_t>, + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int64_t, int32_t>}; fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type, int64_t, + int64_t>, + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int64_t, int64_t>}; fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type, int64_t, + float>, + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int64_t, float>}; fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_LNG][eft_DBL] = { eft_DBL, (void *)dpnp_copysign_c_ext}; fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type, float, + int32_t>, + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + float, int32_t>}; fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type, float, + int64_t>, + get_default_floating_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + float, int64_t>}; fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_FLT][eft_FLT] = { eft_FLT, (void *)dpnp_copysign_c_ext}; fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_FLT][eft_DBL] = { From 2685ce76de446508f4a90293ba1854077ca24725 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 27 Jul 2023 01:20:40 +0200 Subject: [PATCH 03/20] Update dpnp.fmod --- .../include/dpnp_gen_2arg_3type_tbl.hpp | 6 +-- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 51 +++++++++++++++++-- 2 files changed, 50 insertions(+), 7 deletions(-) diff --git a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp index fac16393156..0065ac3e443 100644 --- a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp +++ b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp @@ -138,9 +138,9 @@ MACRO_2ARG_3TYPES_OP(dpnp_divide_c, std::complex)) MACRO_2ARG_3TYPES_OP(dpnp_fmod_c, - sycl::fmod((double)input1_elem, (double)input2_elem), - nullptr, - std::false_type, + dispatch_fmod_op(input1_elem, input2_elem), + x1 % x2, + MACRO_UNPACK_TYPES(std::int32_t, std::int64_t), oneapi::mkl::vm::fmod, MACRO_UNPACK_TYPES(float, double)) diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 73e861ef4d5..7ac41c78309 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -819,6 +819,17 @@ constexpr T dispatch_sign_op(T elem) } } +template +constexpr auto dispatch_fmod_op(T elem1, T elem2) +{ + if constexpr (is_any_v) { + return elem1 % elem2; + } + else { + return sycl::fmod(elem1, elem2); + } +} + #define MACRO_1ARG_1TYPE_OP(__name__, __operation1__, __operation2__) \ template \ class __name__##_kernel; \ @@ -2016,7 +2027,15 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_INT][eft_LNG] = { eft_LNG, (void *)dpnp_fmod_c_ext}; fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_fmod_c_ext}; + get_default_floating_type(), + (void *)dpnp_fmod_c_ext< + func_type_map_t::find_type, int32_t, + float>, + get_default_floating_type(), + (void *) + dpnp_fmod_c_ext()>, + int32_t, float>}; fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_INT][eft_DBL] = { eft_DBL, (void *)dpnp_fmod_c_ext}; fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_LNG][eft_INT] = { @@ -2024,13 +2043,37 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_LNG][eft_LNG] = { eft_LNG, (void *)dpnp_fmod_c_ext}; fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_fmod_c_ext}; + get_default_floating_type(), + (void *)dpnp_fmod_c_ext< + func_type_map_t::find_type, int64_t, + float>, + get_default_floating_type(), + (void *) + dpnp_fmod_c_ext()>, + int64_t, float>}; fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_LNG][eft_DBL] = { eft_DBL, (void *)dpnp_fmod_c_ext}; fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_fmod_c_ext}; + get_default_floating_type(), + (void *)dpnp_fmod_c_ext< + func_type_map_t::find_type, float, + int32_t>, + get_default_floating_type(), + (void *) + dpnp_fmod_c_ext()>, + float, int32_t>}; fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_fmod_c_ext}; + get_default_floating_type(), + (void *)dpnp_fmod_c_ext< + func_type_map_t::find_type, float, + int64_t>, + get_default_floating_type(), + (void *) + dpnp_fmod_c_ext()>, + float, int64_t>}; fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_FLT][eft_FLT] = { eft_FLT, (void *)dpnp_fmod_c_ext}; fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_FLT][eft_DBL] = { From 4af264e7d1ee4f03b0a0108bfae75f81d7a78269 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 27 Jul 2023 11:25:22 +0200 Subject: [PATCH 04/20] --amend --- tests/test_mathematical.py | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index 873a134a9c1..3d5adfe289e 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -162,6 +162,16 @@ def test_divide(self, dtype, lhs, rhs): "dtype", get_all_dtypes(no_bool=True, no_complex=True) ) def test_fmod(self, dtype, lhs, rhs): + if dtype == dpnp.float32 and rhs == 0.3: + """ + Due to some reason NumPy behaves incorrectly, when: + >>> numpy.fmod(numpy.array([3.9], dtype=numpy.float32), 0.3) + array([0.29999995], dtype=float32) + while dpnp returns something around zero which is expected: + >>> dpnp.fmod(dpnp.array([3.9], dtype=dpnp.float32), 0.3) + array([9.53674318e-08]) + """ + pytest.skip("missaligned with numpy results") self._test_mathematical("fmod", dtype, lhs, rhs) @pytest.mark.parametrize("dtype", get_all_dtypes(no_complex=True)) From f901581da9519e08528f1816736e0552b426dc7f Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 27 Jul 2023 01:26:43 +0200 Subject: [PATCH 05/20] Update dpnp.hypot --- .../include/dpnp_gen_2arg_3type_tbl.hpp | 6 +- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 80 +++++++++++++++++-- 2 files changed, 75 insertions(+), 11 deletions(-) diff --git a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp index 0065ac3e443..8b26f4b56e7 100644 --- a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp +++ b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp @@ -145,9 +145,9 @@ MACRO_2ARG_3TYPES_OP(dpnp_fmod_c, MACRO_UNPACK_TYPES(float, double)) MACRO_2ARG_3TYPES_OP(dpnp_hypot_c, - sycl::hypot((double)input1_elem, (double)input2_elem), - nullptr, - std::false_type, + sycl::hypot(input1_elem, input2_elem), + sycl::hypot(x1, x2), + MACRO_UNPACK_TYPES(float, double), oneapi::mkl::vm::hypot, MACRO_UNPACK_TYPES(float, double)) diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 7ac41c78309..b354f4e7d73 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -2121,25 +2121,89 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) eft_DBL, (void *)dpnp_hypot_c_default}; fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_INT][eft_INT] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; + get_default_floating_type(), + (void *)dpnp_hypot_c_ext< + func_type_map_t::find_type, int32_t, + int32_t>, + get_default_floating_type(), + (void *) + dpnp_hypot_c_ext()>, + int32_t, int32_t>}; fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_INT][eft_LNG] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; + get_default_floating_type(), + (void *)dpnp_hypot_c_ext< + func_type_map_t::find_type, int32_t, + int64_t>, + get_default_floating_type(), + (void *) + dpnp_hypot_c_ext()>, + int32_t, int64_t>}; fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; + get_default_floating_type(), + (void *)dpnp_hypot_c_ext< + func_type_map_t::find_type, int32_t, + float>, + get_default_floating_type(), + (void *) + dpnp_hypot_c_ext()>, + int32_t, float>}; fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_INT][eft_DBL] = { eft_DBL, (void *)dpnp_hypot_c_ext}; fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_LNG][eft_INT] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; + get_default_floating_type(), + (void *)dpnp_hypot_c_ext< + func_type_map_t::find_type, int64_t, + int32_t>, + get_default_floating_type(), + (void *) + dpnp_hypot_c_ext()>, + int64_t, int32_t>}; fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_LNG][eft_LNG] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; + get_default_floating_type(), + (void *)dpnp_hypot_c_ext< + func_type_map_t::find_type, int64_t, + int64_t>, + get_default_floating_type(), + (void *) + dpnp_hypot_c_ext()>, + int64_t, int64_t>}; fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; + get_default_floating_type(), + (void *)dpnp_hypot_c_ext< + func_type_map_t::find_type, int64_t, + float>, + get_default_floating_type(), + (void *) + dpnp_hypot_c_ext()>, + int64_t, float>}; fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_LNG][eft_DBL] = { eft_DBL, (void *)dpnp_hypot_c_ext}; fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; + get_default_floating_type(), + (void *)dpnp_hypot_c_ext< + func_type_map_t::find_type, float, + int32_t>, + get_default_floating_type(), + (void *) + dpnp_hypot_c_ext()>, + float, int32_t>}; fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; + get_default_floating_type(), + (void *)dpnp_hypot_c_ext< + func_type_map_t::find_type, float, + int64_t>, + get_default_floating_type(), + (void *) + dpnp_hypot_c_ext()>, + float, int64_t>}; fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_FLT][eft_FLT] = { eft_FLT, (void *)dpnp_hypot_c_ext}; fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_FLT][eft_DBL] = { From fdb13b0c1259c874156525519ed3d40e027ce0db Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 27 Jul 2023 01:33:42 +0200 Subject: [PATCH 06/20] Update dpnp.maximum --- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 40 ++++++++++++++++++--- 1 file changed, 36 insertions(+), 4 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index b354f4e7d73..cadfa6ed345 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -2255,7 +2255,15 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_INT][eft_LNG] = { eft_LNG, (void *)dpnp_maximum_c_ext}; fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_maximum_c_ext}; + get_default_floating_type(), + (void *)dpnp_maximum_c_ext< + func_type_map_t::find_type, int32_t, + float>, + get_default_floating_type(), + (void *)dpnp_maximum_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int32_t, float>}; fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_INT][eft_DBL] = { eft_DBL, (void *)dpnp_maximum_c_ext}; fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_LNG][eft_INT] = { @@ -2263,13 +2271,37 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_LNG][eft_LNG] = { eft_LNG, (void *)dpnp_maximum_c_ext}; fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_maximum_c_ext}; + get_default_floating_type(), + (void *)dpnp_maximum_c_ext< + func_type_map_t::find_type, int64_t, + float>, + get_default_floating_type(), + (void *)dpnp_maximum_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int64_t, float>}; fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_LNG][eft_DBL] = { eft_DBL, (void *)dpnp_maximum_c_ext}; fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_maximum_c_ext}; + get_default_floating_type(), + (void *)dpnp_maximum_c_ext< + func_type_map_t::find_type, float, + int32_t>, + get_default_floating_type(), + (void *)dpnp_maximum_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + float, int32_t>}; fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_maximum_c_ext}; + get_default_floating_type(), + (void *)dpnp_maximum_c_ext< + func_type_map_t::find_type, float, + int64_t>, + get_default_floating_type(), + (void *)dpnp_maximum_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + float, int64_t>}; fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_FLT][eft_FLT] = { eft_FLT, (void *)dpnp_maximum_c_ext}; fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_FLT][eft_DBL] = { From 01a2fc2e2df781719e8368bf9a36f2d101fbdc8a Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 27 Jul 2023 01:40:23 +0200 Subject: [PATCH 07/20] Update dpnp.minimum --- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 40 ++++++++++++++++++--- 1 file changed, 36 insertions(+), 4 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index cadfa6ed345..100ecd20c74 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -2353,7 +2353,15 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_INT][eft_LNG] = { eft_LNG, (void *)dpnp_minimum_c_ext}; fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_minimum_c_ext}; + get_default_floating_type(), + (void *)dpnp_minimum_c_ext< + func_type_map_t::find_type, int32_t, + float>, + get_default_floating_type(), + (void *)dpnp_minimum_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int32_t, float>}; fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_INT][eft_DBL] = { eft_DBL, (void *)dpnp_minimum_c_ext}; fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_LNG][eft_INT] = { @@ -2361,13 +2369,37 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_LNG][eft_LNG] = { eft_LNG, (void *)dpnp_minimum_c_ext}; fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_minimum_c_ext}; + get_default_floating_type(), + (void *)dpnp_minimum_c_ext< + func_type_map_t::find_type, int64_t, + float>, + get_default_floating_type(), + (void *)dpnp_minimum_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + int64_t, float>}; fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_LNG][eft_DBL] = { eft_DBL, (void *)dpnp_minimum_c_ext}; fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_minimum_c_ext}; + get_default_floating_type(), + (void *)dpnp_minimum_c_ext< + func_type_map_t::find_type, float, + int32_t>, + get_default_floating_type(), + (void *)dpnp_minimum_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + float, int32_t>}; fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_minimum_c_ext}; + get_default_floating_type(), + (void *)dpnp_minimum_c_ext< + func_type_map_t::find_type, float, + int64_t>, + get_default_floating_type(), + (void *)dpnp_minimum_c_ext< + func_type_map_t::find_type< + get_default_floating_type()>, + float, int64_t>}; fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_FLT][eft_FLT] = { eft_FLT, (void *)dpnp_minimum_c_ext}; fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_FLT][eft_DBL] = { From d33f785aa4f7914237059b3a0876e43a641ea487 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 27 Jul 2023 12:01:23 +0200 Subject: [PATCH 08/20] Update dpnp.cross --- .../kernels/dpnp_krnl_mathematical.cpp | 40 +++++++++++++++++-- 1 file changed, 36 insertions(+), 4 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp index 5957cb4a699..ebab980d5fb 100644 --- a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp @@ -1223,7 +1223,15 @@ void func_map_init_mathematical(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_INT][eft_LNG] = { eft_LNG, (void *)dpnp_cross_ext_c}; fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_cross_ext_c}; + get_default_floating_type(), + (void *)dpnp_cross_ext_c< + func_type_map_t::find_type, int32_t, + float>, + get_default_floating_type(), + (void *) + dpnp_cross_ext_c()>, + int32_t, float>}; fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_INT][eft_DBL] = { eft_DBL, (void *)dpnp_cross_ext_c}; fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_LNG][eft_INT] = { @@ -1231,13 +1239,37 @@ void func_map_init_mathematical(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_LNG][eft_LNG] = { eft_LNG, (void *)dpnp_cross_ext_c}; fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_cross_ext_c}; + get_default_floating_type(), + (void *)dpnp_cross_ext_c< + func_type_map_t::find_type, int64_t, + float>, + get_default_floating_type(), + (void *) + dpnp_cross_ext_c()>, + int64_t, float>}; fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_LNG][eft_DBL] = { eft_DBL, (void *)dpnp_cross_ext_c}; fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_cross_ext_c}; + get_default_floating_type(), + (void *)dpnp_cross_ext_c< + func_type_map_t::find_type, float, + int32_t>, + get_default_floating_type(), + (void *) + dpnp_cross_ext_c()>, + float, int32_t>}; fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_cross_ext_c}; + get_default_floating_type(), + (void *)dpnp_cross_ext_c< + func_type_map_t::find_type, float, + int64_t>, + get_default_floating_type(), + (void *) + dpnp_cross_ext_c()>, + float, int64_t>}; fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_FLT][eft_FLT] = { eft_FLT, (void *)dpnp_cross_ext_c}; fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_FLT][eft_DBL] = { From cd1a0f2b3ba375d2e2fcdb13c0d03aee07216155 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 27 Jul 2023 13:36:43 +0200 Subject: [PATCH 09/20] Update test_mathematical --- tests/test_mathematical.py | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index 3d5adfe289e..9d70a3053e2 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -13,6 +13,7 @@ from .helper import ( get_all_dtypes, get_float_complex_dtypes, + has_support_aspect64, is_cpu_device, is_win_platform, ) @@ -162,16 +163,16 @@ def test_divide(self, dtype, lhs, rhs): "dtype", get_all_dtypes(no_bool=True, no_complex=True) ) def test_fmod(self, dtype, lhs, rhs): - if dtype == dpnp.float32 and rhs == 0.3: + if dtype == None and rhs == 0.3 and not has_support_aspect64(): """ Due to some reason NumPy behaves incorrectly, when: >>> numpy.fmod(numpy.array([3.9], dtype=numpy.float32), 0.3) array([0.29999995], dtype=float32) - while dpnp returns something around zero which is expected: - >>> dpnp.fmod(dpnp.array([3.9], dtype=dpnp.float32), 0.3) + while numpy with float64 returns something around zero which is expected: + >>> numpy.fmod(numpy.array([3.9], dtype=numpy.float64), 0.3) array([9.53674318e-08]) """ - pytest.skip("missaligned with numpy results") + pytest.skip("missaligned between numpy results") self._test_mathematical("fmod", dtype, lhs, rhs) @pytest.mark.parametrize("dtype", get_all_dtypes(no_complex=True)) @@ -392,7 +393,7 @@ def test_negative(data, dtype): result = dpnp.negative(dpnp_a) expected = numpy.negative(np_a) - assert_array_equal(result, expected) + assert_allclose(result, expected) @pytest.mark.parametrize("val_type", get_all_dtypes(no_none=True)) From 116dbf0ecb0e61511b5f6c9890cdf8747a73e38f Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Wed, 2 Aug 2023 21:08:02 +0200 Subject: [PATCH 10/20] Add a new template function get_res_type --- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 579 +++--------------- .../kernels/dpnp_krnl_mathematical.cpp | 93 +-- dpnp/backend/src/dpnp_fptr.hpp | 56 ++ 3 files changed, 177 insertions(+), 551 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 100ecd20c74..ee836fb4248 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -1655,12 +1655,102 @@ static void func_map_elemwise_2arg_3type_core(func_map_t &fmap) ...); } +template +static void func_map_elemwise_2arg_3type_short_core(func_map_t &fmap) +{ + ((fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][FT1][FTs] = + {get_res_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type()>, + func_type_map_t::find_type, + func_type_map_t::find_type>, + get_res_type(), + (void *)dpnp_arctan2_c_ext()>, + func_type_map_t::find_type, + func_type_map_t::find_type>}), + ...); + ((fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][FT1][FTs] = + {get_res_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type()>, + func_type_map_t::find_type, + func_type_map_t::find_type>, + get_res_type(), + (void *)dpnp_copysign_c_ext()>, + func_type_map_t::find_type, + func_type_map_t::find_type>}), + ...); + ((fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][FT1][FTs] = + {get_res_type(), + (void *) + dpnp_fmod_c_ext()>, + func_type_map_t::find_type, + func_type_map_t::find_type>, + get_res_type(), + (void *)dpnp_fmod_c_ext< + func_type_map_t::find_type< + get_res_type()>, + func_type_map_t::find_type, + func_type_map_t::find_type>}), + ...); + ((fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][FT1][FTs] = + {get_res_type(), + (void *)dpnp_hypot_c_ext< + func_type_map_t::find_type()>, + func_type_map_t::find_type, + func_type_map_t::find_type>, + get_res_type(), + (void *) + dpnp_hypot_c_ext()>, + func_type_map_t::find_type, + func_type_map_t::find_type>}), + ...); + ((fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][FT1][FTs] = + {get_res_type(), + (void *)dpnp_maximum_c_ext< + func_type_map_t::find_type< + get_res_type()>, + func_type_map_t::find_type, + func_type_map_t::find_type>, + get_res_type(), + (void *)dpnp_maximum_c_ext< + func_type_map_t::find_type< + get_res_type()>, + func_type_map_t::find_type, + func_type_map_t::find_type>}), + ...); + ((fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][FT1][FTs] = + {get_res_type(), + (void *)dpnp_minimum_c_ext< + func_type_map_t::find_type< + get_res_type()>, + func_type_map_t::find_type, + func_type_map_t::find_type>, + get_res_type(), + (void *)dpnp_minimum_c_ext< + func_type_map_t::find_type< + get_res_type()>, + func_type_map_t::find_type, + func_type_map_t::find_type>}), + ...); +} + template static void func_map_elemwise_2arg_3type_helper(func_map_t &fmap) { ((func_map_elemwise_2arg_3type_core(fmap)), ...); } +template +static void func_map_elemwise_2arg_3type_short_helper(func_map_t &fmap) +{ + ((func_map_elemwise_2arg_3type_short_core(fmap)), ...); +} + static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) { fmap[DPNPFuncName::DPNP_FN_ADD][eft_INT][eft_INT] = { @@ -1729,103 +1819,6 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_ARCTAN2][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_arctan2_c_default}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_INT][eft_INT] = { - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type, int32_t, - int32_t>, - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int32_t, int32_t>}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_INT][eft_LNG] = { - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type, int32_t, - int64_t>, - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int32_t, int64_t>}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_INT][eft_FLT] = { - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type, int32_t, - float>, - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int32_t, float>}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_LNG][eft_INT] = { - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type, int64_t, - int32_t>, - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int64_t, int32_t>}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_LNG][eft_LNG] = { - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type, int64_t, - int64_t>, - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int64_t, int64_t>}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_LNG][eft_FLT] = { - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type, int64_t, - float>, - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int64_t, float>}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_FLT][eft_INT] = { - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type, float, - int32_t>, - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - float, int32_t>}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_FLT][eft_LNG] = { - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type, float, - int64_t>, - get_default_floating_type(), - (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - float, int64_t>}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_arctan2_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; - fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_arctan2_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_copysign_c_default}; fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_INT][eft_LNG] = { @@ -1859,103 +1852,6 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_COPYSIGN][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_copysign_c_default}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_INT][eft_INT] = { - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type, int32_t, - int32_t>, - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int32_t, int32_t>}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_INT][eft_LNG] = { - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type, int32_t, - int64_t>, - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int32_t, int64_t>}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_INT][eft_FLT] = { - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type, int32_t, - float>, - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int32_t, float>}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_LNG][eft_INT] = { - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type, int64_t, - int32_t>, - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int64_t, int32_t>}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_LNG][eft_LNG] = { - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type, int64_t, - int64_t>, - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int64_t, int64_t>}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_LNG][eft_FLT] = { - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type, int64_t, - float>, - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int64_t, float>}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_FLT][eft_INT] = { - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type, float, - int32_t>, - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - float, int32_t>}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_FLT][eft_LNG] = { - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type, float, - int64_t>, - get_default_floating_type(), - (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - float, int64_t>}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_copysign_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; - fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_copysign_c_ext}; - fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_divide_c_default}; fmap[DPNPFuncName::DPNP_FN_DIVIDE][eft_INT][eft_LNG] = { @@ -2022,71 +1918,6 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_FMOD][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_fmod_c_default}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_fmod_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_fmod_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_INT][eft_FLT] = { - get_default_floating_type(), - (void *)dpnp_fmod_c_ext< - func_type_map_t::find_type, int32_t, - float>, - get_default_floating_type(), - (void *) - dpnp_fmod_c_ext()>, - int32_t, float>}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_fmod_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_fmod_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_fmod_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_LNG][eft_FLT] = { - get_default_floating_type(), - (void *)dpnp_fmod_c_ext< - func_type_map_t::find_type, int64_t, - float>, - get_default_floating_type(), - (void *) - dpnp_fmod_c_ext()>, - int64_t, float>}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_fmod_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_FLT][eft_INT] = { - get_default_floating_type(), - (void *)dpnp_fmod_c_ext< - func_type_map_t::find_type, float, - int32_t>, - get_default_floating_type(), - (void *) - dpnp_fmod_c_ext()>, - float, int32_t>}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_FLT][eft_LNG] = { - get_default_floating_type(), - (void *)dpnp_fmod_c_ext< - func_type_map_t::find_type, float, - int64_t>, - get_default_floating_type(), - (void *) - dpnp_fmod_c_ext()>, - float, int64_t>}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_fmod_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_fmod_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_fmod_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_fmod_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_fmod_c_ext}; - fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_fmod_c_ext}; - fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_hypot_c_default}; fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_INT][eft_LNG] = { @@ -2120,103 +1951,6 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_HYPOT][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_hypot_c_default}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_INT][eft_INT] = { - get_default_floating_type(), - (void *)dpnp_hypot_c_ext< - func_type_map_t::find_type, int32_t, - int32_t>, - get_default_floating_type(), - (void *) - dpnp_hypot_c_ext()>, - int32_t, int32_t>}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_INT][eft_LNG] = { - get_default_floating_type(), - (void *)dpnp_hypot_c_ext< - func_type_map_t::find_type, int32_t, - int64_t>, - get_default_floating_type(), - (void *) - dpnp_hypot_c_ext()>, - int32_t, int64_t>}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_INT][eft_FLT] = { - get_default_floating_type(), - (void *)dpnp_hypot_c_ext< - func_type_map_t::find_type, int32_t, - float>, - get_default_floating_type(), - (void *) - dpnp_hypot_c_ext()>, - int32_t, float>}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_LNG][eft_INT] = { - get_default_floating_type(), - (void *)dpnp_hypot_c_ext< - func_type_map_t::find_type, int64_t, - int32_t>, - get_default_floating_type(), - (void *) - dpnp_hypot_c_ext()>, - int64_t, int32_t>}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_LNG][eft_LNG] = { - get_default_floating_type(), - (void *)dpnp_hypot_c_ext< - func_type_map_t::find_type, int64_t, - int64_t>, - get_default_floating_type(), - (void *) - dpnp_hypot_c_ext()>, - int64_t, int64_t>}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_LNG][eft_FLT] = { - get_default_floating_type(), - (void *)dpnp_hypot_c_ext< - func_type_map_t::find_type, int64_t, - float>, - get_default_floating_type(), - (void *) - dpnp_hypot_c_ext()>, - int64_t, float>}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_FLT][eft_INT] = { - get_default_floating_type(), - (void *)dpnp_hypot_c_ext< - func_type_map_t::find_type, float, - int32_t>, - get_default_floating_type(), - (void *) - dpnp_hypot_c_ext()>, - float, int32_t>}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_FLT][eft_LNG] = { - get_default_floating_type(), - (void *)dpnp_hypot_c_ext< - func_type_map_t::find_type, float, - int64_t>, - get_default_floating_type(), - (void *) - dpnp_hypot_c_ext()>, - float, int64_t>}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_hypot_c_ext}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; - fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_hypot_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_INT][eft_INT] = { eft_INT, (void *)dpnp_maximum_c_default}; fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_INT][eft_LNG] = { @@ -2250,71 +1984,6 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_MAXIMUM][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_maximum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_maximum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_maximum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_INT][eft_FLT] = { - get_default_floating_type(), - (void *)dpnp_maximum_c_ext< - func_type_map_t::find_type, int32_t, - float>, - get_default_floating_type(), - (void *)dpnp_maximum_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int32_t, float>}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_maximum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_maximum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_maximum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_LNG][eft_FLT] = { - get_default_floating_type(), - (void *)dpnp_maximum_c_ext< - func_type_map_t::find_type, int64_t, - float>, - get_default_floating_type(), - (void *)dpnp_maximum_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int64_t, float>}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_maximum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_FLT][eft_INT] = { - get_default_floating_type(), - (void *)dpnp_maximum_c_ext< - func_type_map_t::find_type, float, - int32_t>, - get_default_floating_type(), - (void *)dpnp_maximum_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - float, int32_t>}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_FLT][eft_LNG] = { - get_default_floating_type(), - (void *)dpnp_maximum_c_ext< - func_type_map_t::find_type, float, - int64_t>, - get_default_floating_type(), - (void *)dpnp_maximum_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - float, int64_t>}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_maximum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_maximum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_maximum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_maximum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_maximum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_maximum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_INT][eft_INT] = { eft_INT, (void *)dpnp_minimum_c_default}; fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_INT][eft_LNG] = { @@ -2348,71 +2017,6 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_MINIMUM][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_minimum_c_default}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_minimum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_minimum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_INT][eft_FLT] = { - get_default_floating_type(), - (void *)dpnp_minimum_c_ext< - func_type_map_t::find_type, int32_t, - float>, - get_default_floating_type(), - (void *)dpnp_minimum_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int32_t, float>}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_minimum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_minimum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_minimum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_LNG][eft_FLT] = { - get_default_floating_type(), - (void *)dpnp_minimum_c_ext< - func_type_map_t::find_type, int64_t, - float>, - get_default_floating_type(), - (void *)dpnp_minimum_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - int64_t, float>}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_minimum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_FLT][eft_INT] = { - get_default_floating_type(), - (void *)dpnp_minimum_c_ext< - func_type_map_t::find_type, float, - int32_t>, - get_default_floating_type(), - (void *)dpnp_minimum_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - float, int32_t>}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_FLT][eft_LNG] = { - get_default_floating_type(), - (void *)dpnp_minimum_c_ext< - func_type_map_t::find_type, float, - int64_t>, - get_default_floating_type(), - (void *)dpnp_minimum_c_ext< - func_type_map_t::find_type< - get_default_floating_type()>, - float, int64_t>}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_minimum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_minimum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_minimum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_minimum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_minimum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_minimum_c_ext}; - fmap[DPNPFuncName::DPNP_FN_MULTIPLY][eft_BLN][eft_BLN] = { eft_BLN, (void *)dpnp_multiply_c_default}; fmap[DPNPFuncName::DPNP_FN_MULTIPLY][eft_BLN][eft_INT] = { @@ -2584,6 +2188,9 @@ static void func_map_init_elemwise_2arg_3type(func_map_t &fmap) func_map_elemwise_2arg_3type_helper(fmap); + func_map_elemwise_2arg_3type_short_helper(fmap); + return; } diff --git a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp index ebab980d5fb..1e40ee04798 100644 --- a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp @@ -1141,6 +1141,31 @@ DPCTLSyclEventRef (*dpnp_trapz_ext_c)(DPCTLSyclQueueRef, const DPCTLEventVectorRef) = dpnp_trapz_c<_DataType_input1, _DataType_input2, _DataType_output>; +template +static void func_map_elemwise_2arg_3type_core(func_map_t &fmap) +{ + ((fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][FT1][FTs] = + {get_res_type(), + (void *)dpnp_cross_ext_c< + func_type_map_t::find_type< + get_res_type()>, + func_type_map_t::find_type, + func_type_map_t::find_type>, + get_res_type(), + (void *)dpnp_cross_ext_c< + func_type_map_t::find_type< + get_res_type()>, + func_type_map_t::find_type, + func_type_map_t::find_type>}), + ...); +} + +template +static void func_map_elemwise_2arg_3type_helper(func_map_t &fmap) +{ + ((func_map_elemwise_2arg_3type_core(fmap)), ...); +} + void func_map_init_mathematical(func_map_t &fmap) { fmap[DPNPFuncName::DPNP_FN_ABSOLUTE][eft_INT][eft_INT] = { @@ -1218,71 +1243,6 @@ void func_map_init_mathematical(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_CROSS][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_cross_default_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_cross_ext_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_cross_ext_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_INT][eft_FLT] = { - get_default_floating_type(), - (void *)dpnp_cross_ext_c< - func_type_map_t::find_type, int32_t, - float>, - get_default_floating_type(), - (void *) - dpnp_cross_ext_c()>, - int32_t, float>}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_cross_ext_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_cross_ext_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_cross_ext_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_LNG][eft_FLT] = { - get_default_floating_type(), - (void *)dpnp_cross_ext_c< - func_type_map_t::find_type, int64_t, - float>, - get_default_floating_type(), - (void *) - dpnp_cross_ext_c()>, - int64_t, float>}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_cross_ext_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_FLT][eft_INT] = { - get_default_floating_type(), - (void *)dpnp_cross_ext_c< - func_type_map_t::find_type, float, - int32_t>, - get_default_floating_type(), - (void *) - dpnp_cross_ext_c()>, - float, int32_t>}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_FLT][eft_LNG] = { - get_default_floating_type(), - (void *)dpnp_cross_ext_c< - func_type_map_t::find_type, float, - int64_t>, - get_default_floating_type(), - (void *) - dpnp_cross_ext_c()>, - float, int64_t>}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_cross_ext_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_cross_ext_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_cross_ext_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_cross_ext_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_cross_ext_c}; - fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_cross_ext_c}; - fmap[DPNPFuncName::DPNP_FN_CUMPROD][eft_INT][eft_INT] = { eft_LNG, (void *)dpnp_cumprod_default_c}; fmap[DPNPFuncName::DPNP_FN_CUMPROD][eft_LNG][eft_LNG] = { @@ -1557,5 +1517,8 @@ void func_map_init_mathematical(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_TRAPZ_EXT][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_trapz_ext_c}; + func_map_elemwise_2arg_3type_helper( + fmap); + return; } diff --git a/dpnp/backend/src/dpnp_fptr.hpp b/dpnp/backend/src/dpnp_fptr.hpp index 446083204fc..ff6f301a804 100644 --- a/dpnp/backend/src/dpnp_fptr.hpp +++ b/dpnp/backend/src/dpnp_fptr.hpp @@ -271,6 +271,62 @@ static constexpr DPNPFuncType get_default_floating_type() : DPNPFuncType::DPNP_FT_FLOAT; } +/** + * A template function that determines the resulting floating-point type + * based on the value of the template parameter has_fp64. + */ +template +static constexpr DPNPFuncType get_res_type() +{ + constexpr auto widest_type = populate_func_types(); + constexpr auto shortes_type = (widest_type == FT1) ? FT2 : FT1; + + // Return integer result type if save_int is True + if constexpr (keep_int::value) { + if constexpr (widest_type == DPNPFuncType::DPNP_FT_INT || + widest_type == DPNPFuncType::DPNP_FT_LONG) + { + return widest_type; + } + } + + // Check for double + if constexpr (widest_type == DPNPFuncType::DPNP_FT_DOUBLE) { + return widest_type; + } + + // Check for float + else if constexpr (widest_type == DPNPFuncType::DPNP_FT_FLOAT) { + // Check if the shortest type is also float + if constexpr (shortes_type == DPNPFuncType::DPNP_FT_FLOAT) { + return widest_type; + } + // Check if the shortest type is int or long + else if constexpr (shortes_type == DPNPFuncType::DPNP_FT_INT || + shortes_type == DPNPFuncType::DPNP_FT_LONG) + { + // Check if has_fp64 + if constexpr (has_fp64::value) { + return DPNPFuncType::DPNP_FT_DOUBLE; + } + else { + return DPNPFuncType::DPNP_FT_FLOAT; + } + } + } + + // Default case + if constexpr (has_fp64::value) { + return DPNPFuncType::DPNP_FT_DOUBLE; + } + else { + return DPNPFuncType::DPNP_FT_FLOAT; + } +} + /** * FPTR interface initialization functions */ From 6252bf4deceae83cf18a02a6b7a2969053932f47 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Wed, 2 Aug 2023 21:15:09 +0200 Subject: [PATCH 11/20] Skip test_remainder on Iris Xe --- tests/test_mathematical.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index 9d70a3053e2..bbe3454b930 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -208,6 +208,9 @@ def test_minimum(self, dtype, lhs, rhs): def test_multiply(self, dtype, lhs, rhs): self._test_mathematical("multiply", dtype, lhs, rhs) + @pytest.mark.skipif( + not has_support_aspect64(), reason="Aborted on Iris Xe: SAT-6039" + ) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @pytest.mark.parametrize( "dtype", get_all_dtypes(no_bool=True, no_complex=True) From d8b664f0bf6487c796dba5054d36f6754f1114f0 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 3 Aug 2023 07:44:41 +0200 Subject: [PATCH 12/20] Fix TestAdd --- tests/test_mathematical.py | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index bbe3454b930..d8811a03a07 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -782,9 +782,9 @@ def test_inplace_strided_out(self, dtype): "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) def test_invalid_shape(self, shape): - dp_array1 = dpnp.arange(10, dtype=dpnp.float64) - dp_array2 = dpnp.arange(5, 15, dtype=dpnp.float64) - dp_out = dpnp.empty(shape, dtype=dpnp.float64) + dp_array1 = dpnp.arange(10, dtype=dpnp.float32) + dp_array2 = dpnp.arange(5, 15, dtype=dpnp.float32) + dp_out = dpnp.empty(shape, dtype=dpnp.float32) with pytest.raises(TypeError): dpnp.add(dp_array1, dp_array2, out=dp_out) From afa61b56eb855bb8d718b662c31e6826cb0c69c9 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 3 Aug 2023 07:45:37 +0200 Subject: [PATCH 13/20] Fix TestMultiply --- tests/test_mathematical.py | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index d8811a03a07..f0cd49b8f5f 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -872,9 +872,9 @@ def test_inplace_strided_out(self, dtype): "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) def test_invalid_shape(self, shape): - dp_array1 = dpnp.arange(10, dtype=dpnp.float64) - dp_array2 = dpnp.arange(5, 15, dtype=dpnp.float64) - dp_out = dpnp.empty(shape, dtype=dpnp.float64) + dp_array1 = dpnp.arange(10, dtype=dpnp.float32) + dp_array2 = dpnp.arange(5, 15, dtype=dpnp.float32) + dp_out = dpnp.empty(shape, dtype=dpnp.float32) with pytest.raises(TypeError): dpnp.multiply(dp_array1, dp_array2, out=dp_out) From fc1f4dca890487987958a9059a479a3eced0dc26 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 3 Aug 2023 07:49:38 +0200 Subject: [PATCH 14/20] Use dispatch_fmod_op in vector implementation --- dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp index 8b26f4b56e7..5a5bfc4d817 100644 --- a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp +++ b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp @@ -139,7 +139,7 @@ MACRO_2ARG_3TYPES_OP(dpnp_divide_c, MACRO_2ARG_3TYPES_OP(dpnp_fmod_c, dispatch_fmod_op(input1_elem, input2_elem), - x1 % x2, + dispatch_fmod_op(x1, x2), MACRO_UNPACK_TYPES(std::int32_t, std::int64_t), oneapi::mkl::vm::fmod, MACRO_UNPACK_TYPES(float, double)) From bba06f44a42839ac49e12f6be0d60552e5a162a5 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 3 Aug 2023 07:58:24 +0200 Subject: [PATCH 15/20] Remove fmap for dpnp_floor_divide_ext --- .../kernels/dpnp_krnl_mathematical.cpp | 33 ------------------- 1 file changed, 33 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp index 1e40ee04798..8f08aa50c15 100644 --- a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp @@ -1334,39 +1334,6 @@ void func_map_init_mathematical(func_map_t &fmap) fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE][eft_DBL][eft_DBL] = { eft_DBL, (void *)dpnp_floor_divide_default_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_INT][eft_LNG] = { - eft_LNG, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_INT][eft_FLT] = { - eft_DBL, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_INT][eft_DBL] = { - eft_DBL, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_LNG][eft_INT] = { - eft_LNG, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_LNG][eft_FLT] = { - eft_DBL, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_LNG][eft_DBL] = { - eft_DBL, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_FLT][eft_INT] = { - eft_DBL, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_FLT][eft_LNG] = { - eft_DBL, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_FLT][eft_DBL] = { - eft_DBL, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_DBL][eft_INT] = { - eft_DBL, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_DBL][eft_LNG] = { - eft_DBL, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_DBL][eft_FLT] = { - eft_DBL, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_FLOOR_DIVIDE_EXT][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_floor_divide_ext_c}; - fmap[DPNPFuncName::DPNP_FN_MODF][eft_INT][eft_INT] = { eft_DBL, (void *)dpnp_modf_default_c}; fmap[DPNPFuncName::DPNP_FN_MODF][eft_LNG][eft_LNG] = { From dab3a045032dac1911d6000e6b9cc22f78204f8d Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 3 Aug 2023 15:53:32 +0200 Subject: [PATCH 16/20] Use dispatch_fmod_vec for vector impl --- dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp | 13 +++++++------ dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 12 ++++++++++++ 2 files changed, 19 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 5a5bfc4d817..8b2cb23ec97 100644 --- a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp +++ b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp @@ -137,12 +137,13 @@ MACRO_2ARG_3TYPES_OP(dpnp_divide_c, std::complex, std::complex)) -MACRO_2ARG_3TYPES_OP(dpnp_fmod_c, - dispatch_fmod_op(input1_elem, input2_elem), - dispatch_fmod_op(x1, x2), - MACRO_UNPACK_TYPES(std::int32_t, std::int64_t), - oneapi::mkl::vm::fmod, - MACRO_UNPACK_TYPES(float, double)) +MACRO_2ARG_3TYPES_OP( + dpnp_fmod_c, + dispatch_fmod_op(input1_elem, input2_elem), + dispatch_fmod_vec(x1, x2), + MACRO_UNPACK_TYPES(std::int32_t, std::int64_t, float, double), + oneapi::mkl::vm::fmod, + MACRO_UNPACK_TYPES(float, double)) MACRO_2ARG_3TYPES_OP(dpnp_hypot_c, sycl::hypot(input1_elem, input2_elem), diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index ee836fb4248..3760d8ddab7 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -830,6 +830,18 @@ constexpr auto dispatch_fmod_op(T elem1, T elem2) } } +template +constexpr auto dispatch_fmod_vec(vecT x1, vecT x2) +{ + using typeVec = typename vecT::element_type; + if constexpr (is_any_v) { + return x1 % x2; + } + else { + return sycl::fmod(x1, x2); + } +} + #define MACRO_1ARG_1TYPE_OP(__name__, __operation1__, __operation2__) \ template \ class __name__##_kernel; \ From c58da8262e2f80d6f7aacd90bdd37aee77fc78a1 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 3 Aug 2023 16:01:15 +0200 Subject: [PATCH 17/20] Rename and update get_res_type func --- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 80 ++++++++++--------- .../kernels/dpnp_krnl_mathematical.cpp | 12 +-- dpnp/backend/src/dpnp_fptr.hpp | 21 +---- 3 files changed, 49 insertions(+), 64 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 3760d8ddab7..9e81a752be9 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -1671,81 +1671,83 @@ template static void func_map_elemwise_2arg_3type_short_core(func_map_t &fmap) { ((fmap[DPNPFuncName::DPNP_FN_ARCTAN2_EXT][FT1][FTs] = - {get_res_type(), + {get_floating_res_type(), (void *)dpnp_arctan2_c_ext< - func_type_map_t::find_type()>, + func_type_map_t::find_type()>, func_type_map_t::find_type, func_type_map_t::find_type>, - get_res_type(), - (void *)dpnp_arctan2_c_ext()>, - func_type_map_t::find_type, - func_type_map_t::find_type>}), + get_floating_res_type(), + (void *)dpnp_arctan2_c_ext< + func_type_map_t::find_type< + get_floating_res_type()>, + func_type_map_t::find_type, + func_type_map_t::find_type>}), ...); ((fmap[DPNPFuncName::DPNP_FN_COPYSIGN_EXT][FT1][FTs] = - {get_res_type(), + {get_floating_res_type(), (void *)dpnp_copysign_c_ext< - func_type_map_t::find_type()>, + func_type_map_t::find_type()>, func_type_map_t::find_type, func_type_map_t::find_type>, - get_res_type(), - (void *)dpnp_copysign_c_ext()>, - func_type_map_t::find_type, - func_type_map_t::find_type>}), + get_floating_res_type(), + (void *)dpnp_copysign_c_ext< + func_type_map_t::find_type< + get_floating_res_type()>, + func_type_map_t::find_type, + func_type_map_t::find_type>}), ...); ((fmap[DPNPFuncName::DPNP_FN_FMOD_EXT][FT1][FTs] = - {get_res_type(), + {get_floating_res_type(), (void *) - dpnp_fmod_c_ext()>, func_type_map_t::find_type, func_type_map_t::find_type>, - get_res_type(), + get_floating_res_type(), (void *)dpnp_fmod_c_ext< - func_type_map_t::find_type< - get_res_type()>, + func_type_map_t::find_type()>, func_type_map_t::find_type, func_type_map_t::find_type>}), ...); ((fmap[DPNPFuncName::DPNP_FN_HYPOT_EXT][FT1][FTs] = - {get_res_type(), + {get_floating_res_type(), (void *)dpnp_hypot_c_ext< - func_type_map_t::find_type()>, + func_type_map_t::find_type()>, func_type_map_t::find_type, func_type_map_t::find_type>, - get_res_type(), - (void *) - dpnp_hypot_c_ext()>, - func_type_map_t::find_type, - func_type_map_t::find_type>}), + get_floating_res_type(), + (void *)dpnp_hypot_c_ext< + func_type_map_t::find_type< + get_floating_res_type()>, + func_type_map_t::find_type, + func_type_map_t::find_type>}), ...); ((fmap[DPNPFuncName::DPNP_FN_MAXIMUM_EXT][FT1][FTs] = - {get_res_type(), + {get_floating_res_type(), (void *)dpnp_maximum_c_ext< - func_type_map_t::find_type< - get_res_type()>, + func_type_map_t::find_type()>, func_type_map_t::find_type, func_type_map_t::find_type>, - get_res_type(), + get_floating_res_type(), (void *)dpnp_maximum_c_ext< - func_type_map_t::find_type< - get_res_type()>, + func_type_map_t::find_type()>, func_type_map_t::find_type, func_type_map_t::find_type>}), ...); ((fmap[DPNPFuncName::DPNP_FN_MINIMUM_EXT][FT1][FTs] = - {get_res_type(), + {get_floating_res_type(), (void *)dpnp_minimum_c_ext< - func_type_map_t::find_type< - get_res_type()>, + func_type_map_t::find_type()>, func_type_map_t::find_type, func_type_map_t::find_type>, - get_res_type(), + get_floating_res_type(), (void *)dpnp_minimum_c_ext< - func_type_map_t::find_type< - get_res_type()>, + func_type_map_t::find_type()>, func_type_map_t::find_type, func_type_map_t::find_type>}), ...); diff --git a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp index 8f08aa50c15..32be7b92983 100644 --- a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp @@ -1145,16 +1145,16 @@ template static void func_map_elemwise_2arg_3type_core(func_map_t &fmap) { ((fmap[DPNPFuncName::DPNP_FN_CROSS_EXT][FT1][FTs] = - {get_res_type(), + {get_floating_res_type(), (void *)dpnp_cross_ext_c< - func_type_map_t::find_type< - get_res_type()>, + func_type_map_t::find_type()>, func_type_map_t::find_type, func_type_map_t::find_type>, - get_res_type(), + get_floating_res_type(), (void *)dpnp_cross_ext_c< - func_type_map_t::find_type< - get_res_type()>, + func_type_map_t::find_type()>, func_type_map_t::find_type, func_type_map_t::find_type>}), ...); diff --git a/dpnp/backend/src/dpnp_fptr.hpp b/dpnp/backend/src/dpnp_fptr.hpp index ff6f301a804..d485711f8f5 100644 --- a/dpnp/backend/src/dpnp_fptr.hpp +++ b/dpnp/backend/src/dpnp_fptr.hpp @@ -279,7 +279,7 @@ template -static constexpr DPNPFuncType get_res_type() +static constexpr DPNPFuncType get_floating_res_type() { constexpr auto widest_type = populate_func_types(); constexpr auto shortes_type = (widest_type == FT1) ? FT2 : FT1; @@ -304,27 +304,10 @@ static constexpr DPNPFuncType get_res_type() if constexpr (shortes_type == DPNPFuncType::DPNP_FT_FLOAT) { return widest_type; } - // Check if the shortest type is int or long - else if constexpr (shortes_type == DPNPFuncType::DPNP_FT_INT || - shortes_type == DPNPFuncType::DPNP_FT_LONG) - { - // Check if has_fp64 - if constexpr (has_fp64::value) { - return DPNPFuncType::DPNP_FT_DOUBLE; - } - else { - return DPNPFuncType::DPNP_FT_FLOAT; - } - } } // Default case - if constexpr (has_fp64::value) { - return DPNPFuncType::DPNP_FT_DOUBLE; - } - else { - return DPNPFuncType::DPNP_FT_FLOAT; - } + return get_default_floating_type(); } /** From ec40fc8d0e13a7012c0c28e05e0faf887baae3fd Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Thu, 3 Aug 2023 16:15:58 +0200 Subject: [PATCH 18/20] Fix remarks --- tests/test_mathematical.py | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py index f0cd49b8f5f..fed1928e076 100644 --- a/tests/test_mathematical.py +++ b/tests/test_mathematical.py @@ -165,10 +165,10 @@ def test_divide(self, dtype, lhs, rhs): def test_fmod(self, dtype, lhs, rhs): if dtype == None and rhs == 0.3 and not has_support_aspect64(): """ - Due to some reason NumPy behaves incorrectly, when: + Due to accuracy reason NumPy behaves differently, when: >>> numpy.fmod(numpy.array([3.9], dtype=numpy.float32), 0.3) array([0.29999995], dtype=float32) - while numpy with float64 returns something around zero which is expected: + while numpy with float64 returns something around zero which is aligned with dpnp: >>> numpy.fmod(numpy.array([3.9], dtype=numpy.float64), 0.3) array([9.53674318e-08]) """ @@ -782,9 +782,9 @@ def test_inplace_strided_out(self, dtype): "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) def test_invalid_shape(self, shape): - dp_array1 = dpnp.arange(10, dtype=dpnp.float32) - dp_array2 = dpnp.arange(5, 15, dtype=dpnp.float32) - dp_out = dpnp.empty(shape, dtype=dpnp.float32) + dp_array1 = dpnp.arange(10) + dp_array2 = dpnp.arange(5, 15) + dp_out = dpnp.empty(shape) with pytest.raises(TypeError): dpnp.add(dp_array1, dp_array2, out=dp_out) @@ -872,9 +872,9 @@ def test_inplace_strided_out(self, dtype): "shape", [(0,), (15,), (2, 2)], ids=["(0,)", "(15, )", "(2,2)"] ) def test_invalid_shape(self, shape): - dp_array1 = dpnp.arange(10, dtype=dpnp.float32) - dp_array2 = dpnp.arange(5, 15, dtype=dpnp.float32) - dp_out = dpnp.empty(shape, dtype=dpnp.float32) + dp_array1 = dpnp.arange(10) + dp_array2 = dpnp.arange(5, 15) + dp_out = dpnp.empty(shape) with pytest.raises(TypeError): dpnp.multiply(dp_array1, dp_array2, out=dp_out) From cbdb52eaf2c6bbbe4a6321dd59244365ffb23892 Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Fri, 4 Aug 2023 12:31:49 +0200 Subject: [PATCH 19/20] Use a common check in dispatch_fmod_op --- dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp | 2 +- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 14 +------------- 2 files changed, 2 insertions(+), 14 deletions(-) diff --git a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp index 8b2cb23ec97..353d1400320 100644 --- a/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp +++ b/dpnp/backend/include/dpnp_gen_2arg_3type_tbl.hpp @@ -140,7 +140,7 @@ MACRO_2ARG_3TYPES_OP(dpnp_divide_c, MACRO_2ARG_3TYPES_OP( dpnp_fmod_c, dispatch_fmod_op(input1_elem, input2_elem), - dispatch_fmod_vec(x1, x2), + dispatch_fmod_op(x1, x2), MACRO_UNPACK_TYPES(std::int32_t, std::int64_t, float, double), oneapi::mkl::vm::fmod, MACRO_UNPACK_TYPES(float, double)) diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 9e81a752be9..a0698e7873e 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -822,7 +822,7 @@ constexpr T dispatch_sign_op(T elem) template constexpr auto dispatch_fmod_op(T elem1, T elem2) { - if constexpr (is_any_v) { + if constexpr (sycl::detail::is_integral::value) { return elem1 % elem2; } else { @@ -830,18 +830,6 @@ constexpr auto dispatch_fmod_op(T elem1, T elem2) } } -template -constexpr auto dispatch_fmod_vec(vecT x1, vecT x2) -{ - using typeVec = typename vecT::element_type; - if constexpr (is_any_v) { - return x1 % x2; - } - else { - return sycl::fmod(x1, x2); - } -} - #define MACRO_1ARG_1TYPE_OP(__name__, __operation1__, __operation2__) \ template \ class __name__##_kernel; \ From ca12b8ebeb326d965d90f2783a4b2eeeb802da9a Mon Sep 17 00:00:00 2001 From: Vladislav Perevezentsev Date: Fri, 4 Aug 2023 17:38:59 +0200 Subject: [PATCH 20/20] Update dispatch_fmod_op --- dpnp/backend/kernels/dpnp_krnl_elemwise.cpp | 21 ++++++++++++++++++++- 1 file changed, 20 insertions(+), 1 deletion(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index a0698e7873e..1bd0d1922e0 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -823,7 +823,26 @@ template constexpr auto dispatch_fmod_op(T elem1, T elem2) { if constexpr (sycl::detail::is_integral::value) { - return elem1 % elem2; + if constexpr (sycl::detail::is_vec::value) { + T rem; + using ElemT = typename T::element_type; +#pragma unroll + for (size_t i = 0; i < rem.size(); i++) { + if (elem2[i] == ElemT(0)) { + rem[i] = ElemT(0); + } + else { + rem[i] = elem1[i] % elem2[i]; + } + } + return rem; + } + else { + if (elem2 == T(0)) { + return T(0); + } + return elem1 % elem2; + } } else { return sycl::fmod(elem1, elem2);