From 110c9e5491b7f0662e6bcaebed14be880ace7eeb Mon Sep 17 00:00:00 2001 From: Natalia Polina Date: Thu, 18 Jul 2024 11:27:50 -0700 Subject: [PATCH 1/7] Added support CFD for dpnp.isclose and dpnp.allclose functions --- dpnp/backend/CMakeLists.txt | 1 - dpnp/backend/kernels/dpnp_krnl_logic.cpp | 499 ------------------ dpnp/backend/src/dpnp_fptr.hpp | 1 - dpnp/backend/src/dpnp_iface_fptr.cpp | 1 - dpnp/dpnp_algo/CMakeLists.txt | 1 - dpnp/dpnp_algo/dpnp_algo.pxd | 8 - dpnp/dpnp_algo/dpnp_algo.pyx | 1 - dpnp/dpnp_algo/dpnp_algo_logic.pxi | 110 ---- dpnp/dpnp_iface_logic.py | 177 ++++--- tests/test_logic.py | 31 +- .../cupy/logic_tests/test_comparison.py | 3 +- 11 files changed, 130 insertions(+), 703 deletions(-) delete mode 100644 dpnp/backend/kernels/dpnp_krnl_logic.cpp delete mode 100644 dpnp/dpnp_algo/dpnp_algo_logic.pxi diff --git a/dpnp/backend/CMakeLists.txt b/dpnp/backend/CMakeLists.txt index 7ed57fd929d..cef60c7632a 100644 --- a/dpnp/backend/CMakeLists.txt +++ b/dpnp/backend/CMakeLists.txt @@ -29,7 +29,6 @@ set(DPNP_SRC kernels/dpnp_krnl_elemwise.cpp kernels/dpnp_krnl_fft.cpp kernels/dpnp_krnl_indexing.cpp - kernels/dpnp_krnl_logic.cpp kernels/dpnp_krnl_mathematical.cpp kernels/dpnp_krnl_random.cpp kernels/dpnp_krnl_reduction.cpp diff --git a/dpnp/backend/kernels/dpnp_krnl_logic.cpp b/dpnp/backend/kernels/dpnp_krnl_logic.cpp deleted file mode 100644 index c8197d72889..00000000000 --- a/dpnp/backend/kernels/dpnp_krnl_logic.cpp +++ /dev/null @@ -1,499 +0,0 @@ -//***************************************************************************** -// Copyright (c) 2016-2024, Intel Corporation -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// - Redistributions of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// - Redistributions in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -// THE POSSIBILITY OF SUCH DAMAGE. -//***************************************************************************** - -#include - -#include "dpnp_fptr.hpp" -#include "dpnp_iface.hpp" -#include "dpnp_iterator.hpp" -#include "dpnpc_memory_adapter.hpp" -#include "queue_sycl.hpp" - -// dpctl tensor headers -#include "kernels/alignment.hpp" - -using dpctl::tensor::kernels::alignment_utils::is_aligned; -using dpctl::tensor::kernels::alignment_utils::required_alignment; - -template -class dpnp_all_c_kernel; - -template -DPCTLSyclEventRef dpnp_all_c(DPCTLSyclQueueRef q_ref, - const void *array1_in, - void *result1, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - static_assert(std::is_same_v<_ResultType, bool>, - "Boolean result type is required"); - - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!array1_in || !result1) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - const _DataType *array_in = static_cast(array1_in); - bool *result = static_cast(result1); - - auto fill_event = q.fill(result, true, 1); - - if (!size) { - event_ref = reinterpret_cast(&fill_event); - return DPCTLEvent_Copy(event_ref); - } - - constexpr size_t lws = 64; - constexpr size_t vec_sz = 8; - - auto gws_range = - sycl::range<1>(((size + lws * vec_sz - 1) / (lws * vec_sz)) * lws); - auto lws_range = sycl::range<1>(lws); - sycl::nd_range<1> gws(gws_range, lws_range); - - auto kernel_parallel_for_func = [=](sycl::nd_item<1> nd_it) { - auto gr = nd_it.get_sub_group(); - const auto max_gr_size = gr.get_max_local_range()[0]; - const size_t start = - vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + - gr.get_group_id()[0] * max_gr_size); - const size_t end = sycl::min(start + vec_sz * max_gr_size, size); - - // each work-item reduces over "vec_sz" elements in the input array - bool local_reduction = sycl::joint_none_of( - gr, &array_in[start], &array_in[end], - [&](_DataType elem) { return elem == static_cast<_DataType>(0); }); - - if (gr.leader() && (local_reduction == false)) { - result[0] = false; - } - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.depends_on(fill_event); - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - auto event = q.submit(kernel_func); - - event_ref = reinterpret_cast(&event); - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_all_c(const void *array1_in, void *result1, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_all_c<_DataType, _ResultType>( - q_ref, array1_in, result1, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_all_default_c)(const void *, - void *, - const size_t) = dpnp_all_c<_DataType, _ResultType>; - -template -DPCTLSyclEventRef (*dpnp_all_ext_c)(DPCTLSyclQueueRef, - const void *, - void *, - const size_t, - const DPCTLEventVectorRef) = - dpnp_all_c<_DataType, _ResultType>; - -template -class dpnp_allclose_kernel; - -template -static sycl::event dpnp_allclose(sycl::queue &q, - const _DataType1 *array1, - const _DataType2 *array2, - bool *result, - const size_t size, - const _TolType rtol_val, - const _TolType atol_val) -{ - sycl::event fill_event = q.fill(result, true, 1); - if (!size) { - return fill_event; - } - - constexpr size_t lws = 64; - constexpr size_t vec_sz = 8; - - auto gws_range = - sycl::range<1>(((size + lws * vec_sz - 1) / (lws * vec_sz)) * lws); - auto lws_range = sycl::range<1>(lws); - sycl::nd_range<1> gws(gws_range, lws_range); - - auto kernel_parallel_for_func = [=](sycl::nd_item<1> nd_it) { - auto gr = nd_it.get_sub_group(); - const auto max_gr_size = gr.get_max_local_range()[0]; - const auto gr_size = gr.get_local_linear_range(); - const size_t start = - vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + - gr.get_group_linear_id() * max_gr_size); - const size_t end = sycl::min(start + vec_sz * gr_size, size); - - // each work-item iterates over "vec_sz" elements in the input arrays - bool partial = true; - - for (size_t i = start + gr.get_local_linear_id(); i < end; i += gr_size) - { - if constexpr (std::is_floating_point_v<_DataType1> && - std::is_floating_point_v<_DataType2>) - { - if (std::isinf(array1[i]) || std::isinf(array2[i])) { - partial &= (array1[i] == array2[i]); - continue; - } - - // workaround for std::inf which does not work on CPU - // [CMPLRLLVM-51856] - if (array1[i] == std::numeric_limits<_DataType1>::infinity()) { - partial &= (array1[i] == array2[i]); - continue; - } - else if (array1[i] == - -std::numeric_limits<_DataType1>::infinity()) { - partial &= (array1[i] == array2[i]); - continue; - } - else if (array2[i] == - std::numeric_limits<_DataType2>::infinity()) { - partial &= (array1[i] == array2[i]); - continue; - } - else if (array2[i] == - -std::numeric_limits<_DataType2>::infinity()) { - partial &= (array1[i] == array2[i]); - continue; - } - } - - // casting integral to floating type to avoid bad behavior - // on abs(MIN_INT), which leads to undefined result - using _Arr2Type = std::conditional_t, - _TolType, _DataType2>; - _Arr2Type arr2 = static_cast<_Arr2Type>(array2[i]); - - partial &= (std::abs(array1[i] - arr2) <= - (atol_val + rtol_val * std::abs(arr2))); - } - partial = sycl::all_of_group(gr, partial); - - if (gr.leader() && (partial == false)) { - result[0] = false; - } - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.depends_on(fill_event); - cgh.parallel_for< - class dpnp_allclose_kernel<_DataType1, _DataType2, _TolType>>( - gws, kernel_parallel_for_func); - }; - - return q.submit(kernel_func); -} - -template -DPCTLSyclEventRef dpnp_allclose_c(DPCTLSyclQueueRef q_ref, - const void *array1_in, - const void *array2_in, - void *result1, - const size_t size, - double rtol_val, - double atol_val, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - static_assert(std::is_same_v<_ResultType, bool>, - "Boolean result type is required"); - - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!array1_in || !result1) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::event event; - - const _DataType1 *array1 = static_cast(array1_in); - const _DataType2 *array2 = static_cast(array2_in); - bool *result = static_cast(result1); - - if (q.get_device().has(sycl::aspect::fp64)) { - event = - dpnp_allclose(q, array1, array2, result, size, rtol_val, atol_val); - } - else { - float rtol = static_cast(rtol_val); - float atol = static_cast(atol_val); - event = dpnp_allclose(q, array1, array2, result, size, rtol, atol); - } - - event_ref = reinterpret_cast(&event); - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_allclose_c(const void *array1_in, - const void *array2_in, - void *result1, - const size_t size, - double rtol_val, - double atol_val) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = - dpnp_allclose_c<_DataType1, _DataType2, _ResultType>( - q_ref, array1_in, array2_in, result1, size, rtol_val, atol_val, - dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); -} - -template -void (*dpnp_allclose_default_c)(const void *, - const void *, - void *, - const size_t, - double, - double) = - dpnp_allclose_c<_DataType1, _DataType2, _ResultType>; - -template -DPCTLSyclEventRef (*dpnp_allclose_ext_c)(DPCTLSyclQueueRef, - const void *, - const void *, - void *, - const size_t, - double, - double, - const DPCTLEventVectorRef) = - dpnp_allclose_c<_DataType1, _DataType2, _ResultType>; - -template -class dpnp_any_c_kernel; - -template -DPCTLSyclEventRef dpnp_any_c(DPCTLSyclQueueRef q_ref, - const void *array1_in, - void *result1, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref) -{ - static_assert(std::is_same_v<_ResultType, bool>, - "Boolean result type is required"); - - // avoid warning unused variable - (void)dep_event_vec_ref; - - DPCTLSyclEventRef event_ref = nullptr; - - if (!array1_in || !result1) { - return event_ref; - } - - sycl::queue q = *(reinterpret_cast(q_ref)); - - const _DataType *array_in = static_cast(array1_in); - bool *result = static_cast(result1); - - auto fill_event = q.fill(result, false, 1); - - if (!size) { - event_ref = reinterpret_cast(&fill_event); - return DPCTLEvent_Copy(event_ref); - } - - constexpr size_t lws = 64; - constexpr size_t vec_sz = 8; - - auto gws_range = - sycl::range<1>(((size + lws * vec_sz - 1) / (lws * vec_sz)) * lws); - auto lws_range = sycl::range<1>(lws); - sycl::nd_range<1> gws(gws_range, lws_range); - - auto kernel_parallel_for_func = [=](sycl::nd_item<1> nd_it) { - auto gr = nd_it.get_sub_group(); - const auto max_gr_size = gr.get_max_local_range()[0]; - const size_t start = - vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + - gr.get_group_id()[0] * max_gr_size); - const size_t end = sycl::min(start + vec_sz * max_gr_size, size); - - // each work-item reduces over "vec_sz" elements in the input array - bool local_reduction = sycl::joint_any_of( - gr, &array_in[start], &array_in[end], - [&](_DataType elem) { return elem != static_cast<_DataType>(0); }); - - if (gr.leader() && (local_reduction == true)) { - result[0] = true; - } - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.depends_on(fill_event); - cgh.parallel_for>( - gws, kernel_parallel_for_func); - }; - - auto event = q.submit(kernel_func); - - event_ref = reinterpret_cast(&event); - return DPCTLEvent_Copy(event_ref); -} - -template -void dpnp_any_c(const void *array1_in, void *result1, const size_t size) -{ - DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); - DPCTLEventVectorRef dep_event_vec_ref = nullptr; - DPCTLSyclEventRef event_ref = dpnp_any_c<_DataType, _ResultType>( - q_ref, array1_in, result1, size, dep_event_vec_ref); - DPCTLEvent_WaitAndThrow(event_ref); - DPCTLEvent_Delete(event_ref); -} - -template -void (*dpnp_any_default_c)(const void *, - void *, - const size_t) = dpnp_any_c<_DataType, _ResultType>; - -template -DPCTLSyclEventRef (*dpnp_any_ext_c)(DPCTLSyclQueueRef, - const void *, - void *, - const size_t, - const DPCTLEventVectorRef) = - dpnp_any_c<_DataType, _ResultType>; - -void func_map_init_logic(func_map_t &fmap) -{ - fmap[DPNPFuncName::DPNP_FN_ALL][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_all_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALL][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_all_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALL][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_all_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALL][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_all_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALL][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_all_default_c}; - - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_INT][eft_INT] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_LNG][eft_INT] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_FLT][eft_INT] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_DBL][eft_INT] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_INT][eft_LNG] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_LNG][eft_LNG] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_FLT][eft_LNG] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_DBL][eft_LNG] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_INT][eft_FLT] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_LNG][eft_FLT] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_FLT][eft_FLT] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_DBL][eft_FLT] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_INT][eft_DBL] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_LNG][eft_DBL] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_FLT][eft_DBL] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_DBL][eft_DBL] = { - eft_BLN, (void *)dpnp_allclose_default_c}; - - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_INT][eft_INT] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_LNG][eft_INT] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_FLT][eft_INT] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_DBL][eft_INT] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_INT][eft_LNG] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_LNG][eft_LNG] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_FLT][eft_LNG] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_DBL][eft_LNG] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_INT][eft_FLT] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_LNG][eft_FLT] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_FLT][eft_FLT] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_DBL][eft_FLT] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_INT][eft_DBL] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_LNG][eft_DBL] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_FLT][eft_DBL] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - fmap[DPNPFuncName::DPNP_FN_ALLCLOSE_EXT][eft_DBL][eft_DBL] = { - eft_BLN, (void *)dpnp_allclose_ext_c}; - - fmap[DPNPFuncName::DPNP_FN_ANY][eft_BLN][eft_BLN] = { - eft_BLN, (void *)dpnp_any_default_c}; - fmap[DPNPFuncName::DPNP_FN_ANY][eft_INT][eft_INT] = { - eft_INT, (void *)dpnp_any_default_c}; - fmap[DPNPFuncName::DPNP_FN_ANY][eft_LNG][eft_LNG] = { - eft_LNG, (void *)dpnp_any_default_c}; - fmap[DPNPFuncName::DPNP_FN_ANY][eft_FLT][eft_FLT] = { - eft_FLT, (void *)dpnp_any_default_c}; - fmap[DPNPFuncName::DPNP_FN_ANY][eft_DBL][eft_DBL] = { - eft_DBL, (void *)dpnp_any_default_c}; - - return; -} diff --git a/dpnp/backend/src/dpnp_fptr.hpp b/dpnp/backend/src/dpnp_fptr.hpp index 5e07b11542d..212ce848a5c 100644 --- a/dpnp/backend/src/dpnp_fptr.hpp +++ b/dpnp/backend/src/dpnp_fptr.hpp @@ -330,7 +330,6 @@ void func_map_init_elemwise(func_map_t &fmap); void func_map_init_fft_func(func_map_t &fmap); void func_map_init_indexing_func(func_map_t &fmap); void func_map_init_linalg(func_map_t &fmap); -void func_map_init_logic(func_map_t &fmap); void func_map_init_mathematical(func_map_t &fmap); void func_map_init_random(func_map_t &fmap); void func_map_init_reduction(func_map_t &fmap); diff --git a/dpnp/backend/src/dpnp_iface_fptr.cpp b/dpnp/backend/src/dpnp_iface_fptr.cpp index f80c5b35863..e2659511d9c 100644 --- a/dpnp/backend/src/dpnp_iface_fptr.cpp +++ b/dpnp/backend/src/dpnp_iface_fptr.cpp @@ -132,7 +132,6 @@ static func_map_t func_map_init() func_map_init_fft_func(fmap); func_map_init_indexing_func(fmap); func_map_init_linalg(fmap); - func_map_init_logic(fmap); func_map_init_mathematical(fmap); func_map_init_random(fmap); func_map_init_reduction(fmap); diff --git a/dpnp/dpnp_algo/CMakeLists.txt b/dpnp/dpnp_algo/CMakeLists.txt index 1aea452c5d9..3711cc744c0 100644 --- a/dpnp/dpnp_algo/CMakeLists.txt +++ b/dpnp/dpnp_algo/CMakeLists.txt @@ -5,7 +5,6 @@ set(dpnp_algo_pyx_deps ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_sorting.pxi ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_mathematical.pxi ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_indexing.pxi - ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_logic.pxi ${CMAKE_CURRENT_SOURCE_DIR}/dpnp_algo_special.pxi ) diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 3b5b2383226..abaee1c09d4 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -33,7 +33,6 @@ from dpnp.dpnp_utils.dpnp_algo_utils cimport dpnp_descriptor cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this namespace for Enum import cdef enum DPNPFuncName "DPNPFuncName": - DPNP_FN_ALLCLOSE_EXT DPNP_FN_CHOOSE_EXT DPNP_FN_CORRELATE_EXT DPNP_FN_DEGREES_EXT @@ -161,13 +160,6 @@ cdef DPNPFuncType dpnp_dtype_to_DPNPFuncType(dtype) cdef dpnp_DPNPFuncType_to_dtype(size_t type) -""" -Logic functions -""" -cpdef dpnp_descriptor dpnp_isclose(dpnp_descriptor input1, dpnp_descriptor input2, - double rtol=*, double atol=*, cpp_bool equal_nan=*) - - """ Trigonometric functions """ diff --git a/dpnp/dpnp_algo/dpnp_algo.pyx b/dpnp/dpnp_algo/dpnp_algo.pyx index d304f1d32d3..ac2ed08e136 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pyx +++ b/dpnp/dpnp_algo/dpnp_algo.pyx @@ -59,7 +59,6 @@ __all__ = [ include "dpnp_algo_indexing.pxi" -include "dpnp_algo_logic.pxi" include "dpnp_algo_mathematical.pxi" include "dpnp_algo_sorting.pxi" include "dpnp_algo_special.pxi" diff --git a/dpnp/dpnp_algo/dpnp_algo_logic.pxi b/dpnp/dpnp_algo/dpnp_algo_logic.pxi deleted file mode 100644 index 8810bfac6b1..00000000000 --- a/dpnp/dpnp_algo/dpnp_algo_logic.pxi +++ /dev/null @@ -1,110 +0,0 @@ -# cython: language_level=3 -# cython: linetrace=True -# -*- coding: utf-8 -*- -# ***************************************************************************** -# Copyright (c) 2016-2024, Intel Corporation -# All rights reserved. -# -# Redistribution and use in source and binary forms, with or without -# modification, are permitted provided that the following conditions are met: -# - Redistributions of source code must retain the above copyright notice, -# this list of conditions and the following disclaimer. -# - Redistributions in binary form must reproduce the above copyright notice, -# this list of conditions and the following disclaimer in the documentation -# and/or other materials provided with the distribution. -# -# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" -# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE -# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE -# ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE -# LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR -# CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF -# SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS -# INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN -# CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) -# ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF -# THE POSSIBILITY OF SUCH DAMAGE. -# ***************************************************************************** - -"""Module Backend (Logic part) - -This module contains interface functions between C backend layer -and the rest of the library - -""" - -# NO IMPORTs here. All imports must be placed into main "dpnp_algo.pyx" file - -__all__ += [ - "dpnp_allclose", - "dpnp_isclose", -] - - -ctypedef c_dpctl.DPCTLSyclEventRef(*custom_allclose_1in_1out_func_ptr_t)(c_dpctl.DPCTLSyclQueueRef, - void * , - void * , - void * , - const size_t, - double, - double, - const c_dpctl.DPCTLEventVectorRef) - - -cpdef utils.dpnp_descriptor dpnp_allclose(utils.dpnp_descriptor array1, - utils.dpnp_descriptor array2, - double rtol_val, - double atol_val): - result_sycl_device, result_usm_type, result_sycl_queue = utils.get_common_usm_allocation(array1, array2) - - cdef utils.dpnp_descriptor result = utils_py.create_output_descriptor_py((1,), - dpnp.bool, - None, - device=result_sycl_device, - usm_type=result_usm_type, - sycl_queue=result_sycl_queue) - - cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(array1.dtype) - cdef DPNPFuncType param2_type = dpnp_dtype_to_DPNPFuncType(array2.dtype) - - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_ALLCLOSE_EXT, param1_type, param2_type) - - result_sycl_queue = result.get_array().sycl_queue - - cdef c_dpctl.SyclQueue q = result_sycl_queue - cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() - - cdef custom_allclose_1in_1out_func_ptr_t func = kernel_data.ptr - - cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, - array1.get_data(), - array2.get_data(), - result.get_data(), - array1.size, - rtol_val, - atol_val, - NULL) # dep_events_ref - - with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) - c_dpctl.DPCTLEvent_Delete(event_ref) - - return result - - -cpdef utils.dpnp_descriptor dpnp_isclose(utils.dpnp_descriptor input1, - utils.dpnp_descriptor input2, - double rtol=1e-05, - double atol=1e-08, - cpp_bool equal_nan=False): - result_sycl_device, result_usm_type, result_sycl_queue = utils.get_common_usm_allocation(input1, input2) - cdef utils.dpnp_descriptor result = utils_py.create_output_descriptor_py(input1.shape, - dpnp.bool, - None, - device=result_sycl_device, - usm_type=result_usm_type, - sycl_queue=result_sycl_queue) - - for i in range(result.size): - result.get_pyobj()[i] = numpy.isclose(input1.get_pyobj()[i], input2.get_pyobj()[i], rtol, atol, equal_nan) - - return result diff --git a/dpnp/dpnp_iface_logic.py b/dpnp/dpnp_iface_logic.py index 96295c314e9..630d5274265 100644 --- a/dpnp/dpnp_iface_logic.py +++ b/dpnp/dpnp_iface_logic.py @@ -50,10 +50,8 @@ import numpy import dpnp -from dpnp.dpnp_algo import dpnp_allclose from dpnp.dpnp_algo.dpnp_elementwise_common import DPNPBinaryFunc, DPNPUnaryFunc from dpnp.dpnp_array import dpnp_array -from dpnp.dpnp_utils import call_origin __all__ = [ "all", @@ -178,7 +176,7 @@ def all(a, /, axis=None, out=None, keepdims=False, *, where=True): return result -def allclose(a, b, rtol=1.0e-5, atol=1.0e-8, **kwargs): +def allclose(a, b, rtol=1.0e-5, atol=1.0e-8, equal_nan=False): """ Returns ``True`` if two arrays are element-wise equal within a tolerance. @@ -193,23 +191,28 @@ def allclose(a, b, rtol=1.0e-5, atol=1.0e-8, **kwargs): For full documentation refer to :obj:`numpy.allclose`. + Parameters + ---------- + a : {dpnp.ndarray, usm_ndarray, scalar} + First input array, expected to have numeric data type. + Both inputs `a` and `b` can not be scalars at the same time. + b : {dpnp.ndarray, usm_ndarray, scalar} + Second input array, also expected to have numeric data type. + Both inputs `a` and `b` can not be scalars at the same time. + rtol : {dpnp.ndarray, usm_ndarray, scalar}, optional + The relative tolerance parameter. Default: ``1e-05``. + atol : {dpnp.ndarray, usm_ndarray, scalar}, optional + The absolute tolerance parameter. Default: ``1e-08``. + equal_nan : bool + Whether to compare ``NaNs`` as equal. If ``True``, ``NaNs`` in `a` will + be considered equal to ``NaNs`` in `b` in the output array. + Returns ------- out : dpnp.ndarray A 0-dim array with ``True`` value if the two arrays are equal within the given tolerance; with ``False`` otherwise. - Limitations - ----------- - Parameters `a` and `b` are supported either as :class:`dpnp.ndarray`, - :class:`dpctl.tensor.usm_ndarray` or scalars, but both `a` and `b` - can not be scalars at the same time. - Keyword argument `kwargs` is currently unsupported. - Otherwise the functions will be executed sequentially on CPU. - Parameters `rtol` and `atol` are supported as scalars. Otherwise - ``TypeError`` exception will be raised. - Input array data types are limited by supported integer and - floating DPNP :ref:`Data types`. See Also -------- @@ -230,6 +233,9 @@ def allclose(a, b, rtol=1.0e-5, atol=1.0e-8, **kwargs): >>> b = np.array([1.0, np.nan]) >>> np.allclose(a, b) array([False]) + >>> np.allclose(a, b, equal_nan=True) + array([ True]) + >>> a = np.array([1.0, np.inf]) >>> b = np.array([1.0, np.inf]) @@ -238,39 +244,7 @@ def allclose(a, b, rtol=1.0e-5, atol=1.0e-8, **kwargs): """ - if isscalar(a) and isscalar(b): - # at least one of inputs has to be an array - pass - elif not ( - dpnp.is_supported_array_or_scalar(a) - and dpnp.is_supported_array_or_scalar(b) - ): - pass - elif kwargs: - pass - else: - if not isscalar(rtol): - raise TypeError( - f"An argument `rtol` must be a scalar, but got {rtol}" - ) - if not isscalar(atol): - raise TypeError( - f"An argument `atol` must be a scalar, but got {atol}" - ) - - if isscalar(a): - a = dpnp.full_like(b, fill_value=a) - elif isscalar(b): - b = dpnp.full_like(a, fill_value=b) - elif a.shape != b.shape: - a, b = dpt.broadcast_arrays(a.get_array(), b.get_array()) - - a_desc = dpnp.get_dpnp_descriptor(a, copy_when_nondefault_queue=False) - b_desc = dpnp.get_dpnp_descriptor(b, copy_when_nondefault_queue=False) - if a_desc and b_desc: - return dpnp_allclose(a_desc, b_desc, rtol, atol).get_pyobj() - - return call_origin(numpy.allclose, a, b, rtol=rtol, atol=atol, **kwargs) + return all(isclose(a, b, rtol=rtol, atol=atol, equal_nan=equal_nan)) def any(a, /, axis=None, out=None, keepdims=False, *, where=True): @@ -572,48 +546,109 @@ def any(a, /, axis=None, out=None, keepdims=False, *, where=True): ) -def isclose(x1, x2, rtol=1e-05, atol=1e-08, equal_nan=False): +def isclose(a, b, rtol=1e-05, atol=1e-08, equal_nan=False): """ Returns a boolean array where two arrays are element-wise equal within a tolerance. For full documentation refer to :obj:`numpy.isclose`. - Limitations - ----------- - `x2` is supported to be integer if `x1` is :class:`dpnp.ndarray` or - at least either `x1` or `x2` should be as :class:`dpnp.ndarray`. - Otherwise the function will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameters + ---------- + a : {dpnp.ndarray, usm_ndarray, scalar} + First input array, expected to have numeric data type. + Both inputs `a` and `b` can not be scalars at the same time. + b : {dpnp.ndarray, usm_ndarray, scalar} + Second input array, also expected to have numeric data type. + Both inputs `a` and `b` can not be scalars at the same time. + rtol : {dpnp.ndarray, usm_ndarray, scalar}, optional + The relative tolerance parameter. Default: ``1e-05``. + atol : {dpnp.ndarray, usm_ndarray, scalar}, optional + The absolute tolerance parameter. Default: ``1e-08``. + equal_nan : bool + Whether to compare ``NaNs`` as equal. If ``True``, ``NaNs`` in `a` will + be considered equal to ``NaNs`` in `b` in the output array. + + Returns + ------- + out : dpnp.ndarray + Returns a boolean array of where `a` and `b` are equal within the given + tolerance. See Also -------- - :obj:`dpnp.allclose` : Returns True if two arrays are element-wise equal - within a tolerance. + :obj:`dpnp.allclose` : Returns ``True`` if two arrays are element-wise + equal within a tolerance. Examples -------- >>> import dpnp as np - >>> x1 = np.array([1e10,1e-7]) - >>> x2 = np.array([1.00001e10,1e-8]) - >>> out = np.isclose(x1, x2) - >>> [i for i in out] - [True, False] + >>> a = np.array([1e10, 1e-7]) + >>> b = np.array([1.00001e10, 1e-8) + >>> np.isclose(a, b) + array([ True, False]) - """ + >>> a = np.array([1e10, 1e-8]) + >>> b = np.array([1.00001e10, 1e-9]) + >>> np.isclose(a, b) + array([ True, True]) - # x1_desc = dpnp.get_dpnp_descriptor(x1) - # x2_desc = dpnp.get_dpnp_descriptor(x2) - # if x1_desc and x2_desc: - # result_obj = dpnp_isclose( - # x1_desc, x2_desc, rtol, atol, equal_nan - # ).get_pyobj() - # return result_obj + >>> a = np.array([1e10, 1e-8]) + >>> b = np.array([1.0001e10, 1e-9]) + >>> np.isclose(a, b) + array([False, True]) - return call_origin( - numpy.isclose, x1, x2, rtol=rtol, atol=atol, equal_nan=equal_nan + >>> a = np.array([1.0, np.nan]) + >>> b = np.array([1.0, np.nan]) + >>> np.isclose(a, b) + array([ True, False]) + >>> np.isclose(a, b, equal_nan=True) + array([ True, True]) + + >>> a = np.array([0.0, 0.0]) + >>> b = np.array([1e-8, 1e-7]) + >>> np.isclose(a, b) + array([True, False]) + >>> b = np.array([1e-100, 1e-7]) + >>> np.isclose(a, b, atol=0.0) + array([False, False]) + + >>> a = np.array([1e-10, 1e-10]) + >>> b = np.array([1e-20, 0.0]) + >>> np.isclose(a, b) + array([ True, True]) + >>> b = np.array([1e-20, 0.999999e-10]) + >>> np.isclose(a, b atol=0.0) + array([False, True]) + + """ + + dpnp.check_supported_arrays_type(a, b, scalar_type=True) + dpnp.check_supported_arrays_type( + rtol, atol, scalar_type=True, all_scalars=True ) + if dpnp.isscalar(b): + dt = dpnp.result_type( + b, dpnp.default_float_type(sycl_queue=a.sycl_queue) + ) + b = dpnp.asarray( + b, dtype=dt, sycl_queue=a.sycl_queue, usm_type=a.usm_type + ) + elif b.dtype.kind == "i": + dt = dpnp.result_type( + b, dpnp.default_float_type(sycl_queue=b.sycl_queue) + ) + b = dpnp.asarray(b, dtype=dt) + + result = less_equal(dpnp.abs(a - b), atol + rtol * dpnp.abs(b)) & isfinite( + b + ) | (a == b) + + if equal_nan: + result |= isnan(a) & isnan(b) + return result + def iscomplex(x): """ diff --git a/tests/test_logic.py b/tests/test_logic.py index 2b5e68d7d72..92b766c3cf4 100644 --- a/tests/test_logic.py +++ b/tests/test_logic.py @@ -139,25 +139,21 @@ def test_broadcast(self, sh_a, sh_b): assert_allclose(dp_res, np_res) def test_input_as_scalars(self): - with pytest.raises(NotImplementedError): + with pytest.raises(TypeError): dpnp.allclose(1.0, 1.0) @pytest.mark.parametrize("val", [[1.0], (-3, 7), numpy.arange(5)]) def test_wrong_input_arrays(self, val): - with pytest.raises(NotImplementedError): + with pytest.raises(TypeError): dpnp.allclose(val, val) - @pytest.mark.parametrize( - "tol", [[0.001], (1.0e-6,), dpnp.array(1.0e-3), numpy.array([1.0e-5])] - ) + @pytest.mark.parametrize("tol", [[0.001], (1.0e-6,), numpy.array([1.0e-5])]) def test_wrong_tols(self, tol): a = dpnp.ones(10) b = dpnp.ones(10) for kw in [{"rtol": tol}, {"atol": tol}, {"rtol": tol, "atol": tol}]: - with pytest.raises( - TypeError, match=r"An argument .* must be a scalar, but got" - ): + with pytest.raises(TypeError): dpnp.allclose(a, b, **kw) @@ -479,3 +475,22 @@ def test_infinity_sign_errors(func): out = dpnp.empty_like(x, dtype="int32") with pytest.raises(ValueError): getattr(dpnp, func)(x, out=out) + + +@pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True, no_complex=True)) +@pytest.mark.parametrize( + "rtol", [1e-05, dpnp.array(1e-05), dpnp.full(10, 1e-05)] +) +@pytest.mark.parametrize( + "atol", [1e-08, dpnp.array(1e-08), dpnp.full(10, 1e-08)] +) +def test_isclose(dtype, rtol, atol): + a = numpy.random.rand(10) + b = a + numpy.random.rand(10) * 1e-8 + + dpnp_a = dpnp.array(a, dtype=dtype) + dpnp_b = dpnp.array(b, dtype=dtype) + + np_res = numpy.isclose(a, b, 1e-05, 1e-08) + dpnp_res = dpnp.isclose(dpnp_a, dpnp_b, rtol, atol) + assert_allclose(dpnp_res, np_res) diff --git a/tests/third_party/cupy/logic_tests/test_comparison.py b/tests/third_party/cupy/logic_tests/test_comparison.py index eed4c7f9b36..fb79c4e509c 100644 --- a/tests/third_party/cupy/logic_tests/test_comparison.py +++ b/tests/third_party/cupy/logic_tests/test_comparison.py @@ -142,7 +142,6 @@ def test_allclose_infinite(self, xp, dtype): @testing.for_float_dtypes() @testing.numpy_cupy_equal() - @pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_allclose_infinite_equal_nan(self, xp, dtype): nan = float("nan") inf = float("inf") @@ -159,7 +158,6 @@ def test_allclose_array_scalar(self, xp, dtype): return xp.allclose(a, b) -@pytest.mark.usefixtures("allow_fall_back_on_numpy") class TestIsclose(unittest.TestCase): @testing.for_all_dtypes(no_complex=True) @testing.numpy_cupy_array_equal() @@ -204,6 +202,7 @@ def test_is_close_array_scalar(self, xp, dtype): b = xp.dtype(xp.dtype).type(0) return xp.isclose(a, b) + @pytest.mark.skip("support for scalar not implemented") @testing.for_all_dtypes(no_complex=True) def test_is_close_scalar_scalar(self, dtype): # cupy.isclose always returns ndarray From 7aebb773247cff3d3cf034dee219b68e41692efe Mon Sep 17 00:00:00 2001 From: Natalia Polina Date: Tue, 23 Jul 2024 15:39:19 -0700 Subject: [PATCH 2/7] Applied review comments --- dpnp/backend/include/dpnp_iface.hpp | 75 ------------- dpnp/backend/include/dpnp_iface_fptr.hpp | 5 - dpnp/dpnp_iface_logic.py | 52 +++++---- tests/test_sycl_queue.py | 3 +- tests/test_usm_type.py | 3 +- .../cupy/logic_tests/test_comparison.py | 106 ++++++++++++++++-- 6 files changed, 127 insertions(+), 117 deletions(-) diff --git a/dpnp/backend/include/dpnp_iface.hpp b/dpnp/backend/include/dpnp_iface.hpp index 4efea15a38b..d3b342bee30 100644 --- a/dpnp/backend/include/dpnp_iface.hpp +++ b/dpnp/backend/include/dpnp_iface.hpp @@ -100,81 +100,6 @@ INP_DLLEXPORT void dpnp_memory_memcpy_c(DPCTLSyclQueueRef q_ref, size_t size_in_bytes); INP_DLLEXPORT void dpnp_memory_memcpy_c(void *dst, const void *src, size_t size_in_bytes); -/** - * @ingroup BACKEND_API - * @brief Test whether all array elements along a given axis evaluate to True. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array. - * @param [out] result Output array. - * @param [in] size Number of input elements in `array`. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_all_c(DPCTLSyclQueueRef q_ref, - const void *array, - void *result, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_all_c(const void *array, void *result, const size_t size); - -/** - * @ingroup BACKEND_API - * @brief Returns True if two arrays are element-wise equal within a tolerance. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array1_in First input array. - * @param [in] array2_in Second input array. - * @param [out] result1 Output array. - * @param [in] size Number of input elements in `array`. - * @param [in] rtol The relative tolerance parameter. - * @param [in] atol The absolute tolerance parameter. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_allclose_c(DPCTLSyclQueueRef q_ref, - const void *array1_in, - const void *array2_in, - void *result1, - const size_t size, - double rtol, - double atol, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void dpnp_allclose_c(const void *array1_in, - const void *array2_in, - void *result1, - const size_t size, - double rtol, - double atol); - -/** - * @ingroup BACKEND_API - * @brief Test whether any array element along a given axis evaluates to True. - * - * @param [in] q_ref Reference to SYCL queue. - * @param [in] array Input array. - * @param [out] result Output array. - * @param [in] size Number of input elements in `array`. - * @param [in] dep_event_vec_ref Reference to vector of SYCL events. - */ -template -INP_DLLEXPORT DPCTLSyclEventRef - dpnp_any_c(DPCTLSyclQueueRef q_ref, - const void *array, - void *result, - const size_t size, - const DPCTLEventVectorRef dep_event_vec_ref); - -template -INP_DLLEXPORT void - dpnp_any_c(const void *array, void *result, const size_t size); /** * @ingroup BACKEND_API diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index 9f9b7a89143..4f8a3f99485 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -59,11 +59,6 @@ enum class DPNPFuncName : size_t { DPNP_FN_NONE, /**< Very first element of the enumeration */ - DPNP_FN_ALL, /**< Used in numpy.all() impl */ - DPNP_FN_ALLCLOSE, /**< Used in numpy.allclose() impl */ - DPNP_FN_ALLCLOSE_EXT, /**< Used in numpy.allclose() impl, requires extra - parameters */ - DPNP_FN_ANY, /**< Used in numpy.any() impl */ DPNP_FN_ARGMAX, /**< Used in numpy.argmax() impl */ DPNP_FN_ARGMIN, /**< Used in numpy.argmin() impl */ DPNP_FN_ARGSORT, /**< Used in numpy.argsort() impl */ diff --git a/dpnp/dpnp_iface_logic.py b/dpnp/dpnp_iface_logic.py index 630d5274265..a24f2584508 100644 --- a/dpnp/dpnp_iface_logic.py +++ b/dpnp/dpnp_iface_logic.py @@ -185,9 +185,9 @@ def allclose(a, b, rtol=1.0e-5, atol=1.0e-8, equal_nan=False): are added together to compare against the absolute difference between `a` and `b`. - If either array contains one or more ``NaNs``, ``False`` is returned. - ``Infs`` are treated as equal if they are in the same place and of the same - sign in both arrays. + ``NaNs`` are treated as equal if they are in the same place and if + ``equal_nan=True``. ``Infs`` are treated as equal if they are in the same + place and of the same sign in both arrays. For full documentation refer to :obj:`numpy.allclose`. @@ -206,6 +206,7 @@ def allclose(a, b, rtol=1.0e-5, atol=1.0e-8, equal_nan=False): equal_nan : bool Whether to compare ``NaNs`` as equal. If ``True``, ``NaNs`` in `a` will be considered equal to ``NaNs`` in `b` in the output array. + Default: ``False``. Returns ------- @@ -227,20 +228,19 @@ def allclose(a, b, rtol=1.0e-5, atol=1.0e-8, equal_nan=False): >>> a = np.array([1e10, 1e-7]) >>> b = np.array([1.00001e10, 1e-8]) >>> np.allclose(a, b) - array([False]) + array(False) >>> a = np.array([1.0, np.nan]) >>> b = np.array([1.0, np.nan]) >>> np.allclose(a, b) - array([False]) + array(False) >>> np.allclose(a, b, equal_nan=True) - array([ True]) - + array(True) >>> a = np.array([1.0, np.inf]) >>> b = np.array([1.0, np.inf]) >>> np.allclose(a, b) - array([ True]) + array(True) """ @@ -551,6 +551,11 @@ def isclose(a, b, rtol=1e-05, atol=1e-08, equal_nan=False): Returns a boolean array where two arrays are element-wise equal within a tolerance. + The tolerance values are positive, typically very small numbers. The + relative difference (`rtol` * abs(`b`)) and the absolute difference `atol` + are added together to compare against the absolute difference between `a` + and `b`. + For full documentation refer to :obj:`numpy.isclose`. Parameters @@ -568,6 +573,7 @@ def isclose(a, b, rtol=1e-05, atol=1e-08, equal_nan=False): equal_nan : bool Whether to compare ``NaNs`` as equal. If ``True``, ``NaNs`` in `a` will be considered equal to ``NaNs`` in `b` in the output array. + Default: ``False``. Returns ------- @@ -584,7 +590,7 @@ def isclose(a, b, rtol=1e-05, atol=1e-08, equal_nan=False): -------- >>> import dpnp as np >>> a = np.array([1e10, 1e-7]) - >>> b = np.array([1.00001e10, 1e-8) + >>> b = np.array([1.00001e10, 1e-8]) >>> np.isclose(a, b) array([ True, False]) @@ -608,7 +614,7 @@ def isclose(a, b, rtol=1e-05, atol=1e-08, equal_nan=False): >>> a = np.array([0.0, 0.0]) >>> b = np.array([1e-8, 1e-7]) >>> np.isclose(a, b) - array([True, False]) + array([ True, False]) >>> b = np.array([1e-100, 1e-7]) >>> np.isclose(a, b, atol=0.0) array([False, False]) @@ -618,7 +624,7 @@ def isclose(a, b, rtol=1e-05, atol=1e-08, equal_nan=False): >>> np.isclose(a, b) array([ True, True]) >>> b = np.array([1e-20, 0.999999e-10]) - >>> np.isclose(a, b atol=0.0) + >>> np.isclose(a, b, atol=0.0) array([False, True]) """ @@ -628,22 +634,22 @@ def isclose(a, b, rtol=1e-05, atol=1e-08, equal_nan=False): rtol, atol, scalar_type=True, all_scalars=True ) + # make sure b is an inexact type to avoid bad behavior on abs(MIN_INT) if dpnp.isscalar(b): - dt = dpnp.result_type( - b, dpnp.default_float_type(sycl_queue=a.sycl_queue) - ) + dt = dpnp.result_type(a, b, 1.0) b = dpnp.asarray( b, dtype=dt, sycl_queue=a.sycl_queue, usm_type=a.usm_type ) - elif b.dtype.kind == "i": - dt = dpnp.result_type( - b, dpnp.default_float_type(sycl_queue=b.sycl_queue) - ) - b = dpnp.asarray(b, dtype=dt) - - result = less_equal(dpnp.abs(a - b), atol + rtol * dpnp.abs(b)) & isfinite( - b - ) | (a == b) + elif dpnp.issubdtype(b, dpnp.integer): + dt = dpnp.result_type(b, 1.0) + b = dpnp.astype(b, dtype=dt) + + _b = dpnp.abs(b) + _b *= rtol + _b += atol + result = less_equal(dpnp.abs(a - b), _b) + result &= isfinite(b) + result |= a == b if equal_nan: result |= isnan(a) & isnan(b) diff --git a/tests/test_sycl_queue.py b/tests/test_sycl_queue.py index 5b6287dc9f0..1a9c5a348b5 100644 --- a/tests/test_sycl_queue.py +++ b/tests/test_sycl_queue.py @@ -743,8 +743,7 @@ def test_2in_1out(func, data1, data2, device): "equal", "greater", "greater_equal", - # TODO: unblock when dpnp.isclose() is updated - # "isclose", + "isclose", "less", "less_equal", "logical_and", diff --git a/tests/test_usm_type.py b/tests/test_usm_type.py index 6cc8d5edd39..32bd94d0586 100644 --- a/tests/test_usm_type.py +++ b/tests/test_usm_type.py @@ -381,8 +381,7 @@ def test_coerced_usm_types_logic_op_1in(op, usm_type_x): "equal", "greater", "greater_equal", - # TODO: unblock when dpnp.isclose() is updated - # "isclose", + "isclose", "less", "less_equal", "logical_and", diff --git a/tests/third_party/cupy/logic_tests/test_comparison.py b/tests/third_party/cupy/logic_tests/test_comparison.py index fb79c4e509c..668c4679f59 100644 --- a/tests/third_party/cupy/logic_tests/test_comparison.py +++ b/tests/third_party/cupy/logic_tests/test_comparison.py @@ -9,6 +9,7 @@ class TestComparison(unittest.TestCase): + @testing.for_all_dtypes(no_complex=True) @testing.numpy_cupy_allclose(atol=1e-5) def check_binary(self, name, xp, dtype): @@ -36,6 +37,7 @@ def test_equal(self): class TestComparisonOperator(unittest.TestCase): + operators = [ operator.lt, operator.le, @@ -75,6 +77,7 @@ def test_binary_array_pyscalar(self, xp, dtype): class TestArrayEqual(unittest.TestCase): + @testing.for_all_dtypes() @testing.numpy_cupy_equal() def test_array_equal_not_equal(self, xp, dtype): @@ -96,6 +99,29 @@ def test_array_equal_diff_length(self, xp, dtype): b = xp.array([1, 2, 3], dtype=dtype) return xp.array_equal(a, b) + @testing.with_requires("numpy>=1.19") + @testing.for_float_dtypes() + @testing.numpy_cupy_equal() + @pytest.mark.skip("Not supported yet") + def test_array_equal_infinite_equal_nan(self, xp, dtype): + nan = float("nan") + inf = float("inf") + ninf = float("-inf") + a = xp.array([0, nan, inf, ninf], dtype=dtype) + b = xp.array([0, nan, inf, ninf], dtype=dtype) + return xp.array_equal(a, b, equal_nan=True) + + @testing.with_requires("numpy>=1.19") + @testing.for_complex_dtypes() + @testing.numpy_cupy_equal() + @pytest.mark.skip("Not supported yet") + def test_array_equal_complex_equal_nan(self, xp, dtype): + a = xp.array([1 + 2j], dtype=dtype) + b = a.copy() + b.imag = xp.nan + a.real = xp.nan + return xp.array_equal(a, b, equal_nan=True) + @testing.numpy_cupy_equal() def test_array_equal_diff_dtypes_not_equal(self, xp): a = xp.array([0.9e-5, 1.1e-5, 100.5, 10.5]) @@ -115,7 +141,66 @@ def test_array_equal_broadcast_not_allowed(self, xp): return xp.array_equal(a, b) +@pytest.mark.skip("dpnp.array_equiv() is not implemented yet") +class TestArrayEquiv(unittest.TestCase): + + @testing.for_all_dtypes() + @testing.numpy_cupy_equal() + def test_array_equiv_equal_not_equal(self, xp, dtype): + a = xp.array([0, 4, 1, 1], dtype=dtype) + b = xp.array([0, 4, 2, 3], dtype=dtype) + return xp.array_equiv(a, b) + + @testing.for_all_dtypes() + @testing.numpy_cupy_equal() + def test_array_equiv_equal_is_equal(self, xp, dtype): + a = xp.array([0, 4, 1, 1], dtype=dtype) + b = xp.array([0, 4, 1, 1], dtype=dtype) + return xp.array_equiv(a, b) + + @testing.for_all_dtypes() + @testing.numpy_cupy_equal() + def test_array_equiv_diff_length(self, xp, dtype): + a = xp.array([0, 4, 0, 5], dtype=dtype) + b = xp.array([0, 4, 0], dtype=dtype) + return xp.array_equiv(a, b) + + @testing.numpy_cupy_equal() + def test_array_equiv_diff_dtypes_not_equal(self, xp): + a = xp.array([0.9e-5, 1.1e-5, 100.5, 10.5]) + b = xp.array([0, 0, 1000, 1000]) + return xp.array_equiv(a, b) + + @testing.numpy_cupy_equal() + def test_array_equiv_diff_dtypes_is_equal(self, xp): + a = xp.array([0.0, 1.0, 100.0, 10.0]) + b = xp.array([0, 1, 100, 10]) + return xp.array_equiv(a, b) + + @testing.for_all_dtypes() + @testing.numpy_cupy_equal() + def test_array_equiv_broadcast1(self, xp, dtype): + a = xp.array([0, 4], dtype) + b = xp.array([[0, 4], [0, 4]], dtype) + return xp.array_equiv(a, b) + + @testing.for_all_dtypes() + @testing.numpy_cupy_equal() + def test_array_equiv_broadcast2(self, xp, dtype): + a = xp.array([0, 4], dtype=dtype) + b = xp.array([[0, 4], [0, 5]], dtype=dtype) + return xp.array_equiv(a, b) + + @testing.for_all_dtypes() + @testing.numpy_cupy_equal() + def test_array_equiv_non_broadcast(self, xp, dtype): + a = xp.array([0, 4], dtype=dtype) + b = xp.array([[0, 4, 0, 4], [0, 4, 0, 4]], dtype=dtype) + return xp.array_equiv(a, b) + + class TestAllclose(unittest.TestCase): + @testing.for_all_dtypes() @testing.numpy_cupy_equal() def test_allclose_finite(self, xp, dtype): @@ -159,20 +244,21 @@ def test_allclose_array_scalar(self, xp, dtype): class TestIsclose(unittest.TestCase): + @testing.for_all_dtypes(no_complex=True) @testing.numpy_cupy_array_equal() def test_is_close_finite(self, xp, dtype): # In numpy<1.10 this test fails when dtype is bool - a = xp.array([0.9e-5, 1.1e-5, 1000 + 1e-4, 1000 - 1e-4], dtype=dtype) - b = xp.array([0, 0, 1000, 1000], dtype=dtype) + a = xp.array([0.9e-5, 1.1e-5, 1000 + 1e-4, 1000 - 1e-4]).astype(dtype) + b = xp.array([0, 0, 1000, 1000]).astype(dtype) return xp.isclose(a, b) @testing.for_all_dtypes(no_complex=True) @testing.numpy_cupy_array_equal() def test_is_close_min_int(self, xp, dtype): # In numpy<1.10 this test fails when dtype is bool - a = xp.array([0], dtype=dtype) - b = xp.array([numpy.iinfo("i").min], dtype=dtype) + a = xp.array([0]).astype(dtype) + b = xp.array([numpy.iinfo("i").min]).astype(dtype) return xp.isclose(a, b) @testing.for_float_dtypes() @@ -181,8 +267,8 @@ def test_is_close_infinite(self, xp, dtype): nan = float("nan") inf = float("inf") ninf = float("-inf") - a = xp.array([0, nan, nan, 0, inf, ninf], dtype=dtype) - b = xp.array([0, nan, 0, nan, inf, ninf], dtype=dtype) + a = xp.array([0, nan, nan, 0, inf, ninf]).astype(dtype) + b = xp.array([0, nan, 0, nan, inf, ninf]).astype(dtype) return xp.isclose(a, b) @testing.for_float_dtypes() @@ -191,18 +277,18 @@ def test_is_close_infinite_equal_nan(self, xp, dtype): nan = float("nan") inf = float("inf") ninf = float("-inf") - a = xp.array([0, nan, inf, ninf], dtype=dtype) - b = xp.array([0, nan, inf, ninf], dtype=dtype) + a = xp.array([0, nan, inf, ninf]).astype(dtype) + b = xp.array([0, nan, inf, ninf]).astype(dtype) return xp.isclose(a, b, equal_nan=True) @testing.for_all_dtypes(no_complex=True) @testing.numpy_cupy_array_equal() def test_is_close_array_scalar(self, xp, dtype): - a = xp.array([0.9e-5, 1.1e-5], dtype=dtype) + a = xp.array([0.9e-5, 1.1e-5]).astype(dtype) b = xp.dtype(xp.dtype).type(0) return xp.isclose(a, b) - @pytest.mark.skip("support for scalar not implemented") + @pytest.mark.skip(reason="Scalar input is not supported") @testing.for_all_dtypes(no_complex=True) def test_is_close_scalar_scalar(self, dtype): # cupy.isclose always returns ndarray From c1abbf1346a81255fa4139f284da3d99817cfcff Mon Sep 17 00:00:00 2001 From: Natalia Polina Date: Wed, 24 Jul 2024 08:13:59 -0700 Subject: [PATCH 3/7] Update dpnp/dpnp_iface_logic.py Co-authored-by: Anton <100830759+antonwolfy@users.noreply.github.com> --- dpnp/dpnp_iface_logic.py | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/dpnp/dpnp_iface_logic.py b/dpnp/dpnp_iface_logic.py index a24f2584508..434b6c94446 100644 --- a/dpnp/dpnp_iface_logic.py +++ b/dpnp/dpnp_iface_logic.py @@ -556,6 +556,10 @@ def isclose(a, b, rtol=1e-05, atol=1e-08, equal_nan=False): are added together to compare against the absolute difference between `a` and `b`. + ``NaNs`` are treated as equal if they are in the same place and if + ``equal_nan=True``. ``Infs`` are treated as equal if they are in the same + place and of the same sign in both arrays. + For full documentation refer to :obj:`numpy.isclose`. Parameters From e0d3aa3247556f0a12af5d83c6011eb602cb8296 Mon Sep 17 00:00:00 2001 From: Natalia Polina Date: Wed, 24 Jul 2024 08:14:08 -0700 Subject: [PATCH 4/7] Update dpnp/dpnp_iface_logic.py Co-authored-by: Anton <100830759+antonwolfy@users.noreply.github.com> --- dpnp/dpnp_iface_logic.py | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/dpnp/dpnp_iface_logic.py b/dpnp/dpnp_iface_logic.py index 434b6c94446..41a07380de2 100644 --- a/dpnp/dpnp_iface_logic.py +++ b/dpnp/dpnp_iface_logic.py @@ -648,10 +648,15 @@ def isclose(a, b, rtol=1e-05, atol=1e-08, equal_nan=False): dt = dpnp.result_type(b, 1.0) b = dpnp.astype(b, dtype=dt) + # Firstly handle finite values: + # result = absolute(a - b) <= atol + rtol * absolute(b) _b = dpnp.abs(b) _b *= rtol _b += atol result = less_equal(dpnp.abs(a - b), _b) + + # Handle "Inf" values: they are treated as equal if they are in the same place + # and of the same sign in both arrays result &= isfinite(b) result |= a == b From 8a2f016cb8a9581a5de85c4fdaef1f7e6efcec8c Mon Sep 17 00:00:00 2001 From: Natalia Polina Date: Wed, 24 Jul 2024 08:21:17 -0700 Subject: [PATCH 5/7] Fix pre-commit --- dpnp/dpnp_iface_logic.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/dpnp/dpnp_iface_logic.py b/dpnp/dpnp_iface_logic.py index 41a07380de2..ea80c54aaf6 100644 --- a/dpnp/dpnp_iface_logic.py +++ b/dpnp/dpnp_iface_logic.py @@ -655,8 +655,8 @@ def isclose(a, b, rtol=1e-05, atol=1e-08, equal_nan=False): _b += atol result = less_equal(dpnp.abs(a - b), _b) - # Handle "Inf" values: they are treated as equal if they are in the same place - # and of the same sign in both arrays + # Handle "Inf" values: they are treated as equal if they are in the same + # place and of the same sign in both arrays result &= isfinite(b) result |= a == b From ff41a08168725515314db64e02cd146c8e3aeb29 Mon Sep 17 00:00:00 2001 From: Natalia Polina Date: Wed, 24 Jul 2024 15:02:55 -0700 Subject: [PATCH 6/7] fix isclose for mix dtype input --- dpnp/dpnp_iface_logic.py | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/dpnp/dpnp_iface_logic.py b/dpnp/dpnp_iface_logic.py index ea80c54aaf6..a32c2de5307 100644 --- a/dpnp/dpnp_iface_logic.py +++ b/dpnp/dpnp_iface_logic.py @@ -640,17 +640,17 @@ def isclose(a, b, rtol=1e-05, atol=1e-08, equal_nan=False): # make sure b is an inexact type to avoid bad behavior on abs(MIN_INT) if dpnp.isscalar(b): - dt = dpnp.result_type(a, b, 1.0) + dt = dpnp.result_type(a, b, 1.0, rtol, atol) b = dpnp.asarray( b, dtype=dt, sycl_queue=a.sycl_queue, usm_type=a.usm_type ) elif dpnp.issubdtype(b, dpnp.integer): - dt = dpnp.result_type(b, 1.0) + dt = dpnp.result_type(b, 1.0, rtol, atol) b = dpnp.astype(b, dtype=dt) # Firstly handle finite values: # result = absolute(a - b) <= atol + rtol * absolute(b) - _b = dpnp.abs(b) + _b = dpnp.abs(b).astype(dpnp.result_type(b, rtol, atol), copy=False) _b *= rtol _b += atol result = less_equal(dpnp.abs(a - b), _b) From f4634a5a64cd767e3e0544925bd148d7c318d8f9 Mon Sep 17 00:00:00 2001 From: Natalia Polina Date: Thu, 25 Jul 2024 09:09:39 -0700 Subject: [PATCH 7/7] Applied review comments --- dpnp/dpnp_iface_logic.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/dpnp/dpnp_iface_logic.py b/dpnp/dpnp_iface_logic.py index a32c2de5307..2f28b3635ec 100644 --- a/dpnp/dpnp_iface_logic.py +++ b/dpnp/dpnp_iface_logic.py @@ -650,7 +650,8 @@ def isclose(a, b, rtol=1e-05, atol=1e-08, equal_nan=False): # Firstly handle finite values: # result = absolute(a - b) <= atol + rtol * absolute(b) - _b = dpnp.abs(b).astype(dpnp.result_type(b, rtol, atol), copy=False) + dt = dpnp.result_type(b, rtol, atol) + _b = dpnp.abs(b, dtype=dt) _b *= rtol _b += atol result = less_equal(dpnp.abs(a - b), _b)