diff --git a/dpnp/backend/include/dpnp_gen_2arg_2type_tbl.hpp b/dpnp/backend/include/dpnp_gen_2arg_2type_tbl.hpp new file mode 100644 index 00000000000..e76c92b47cd --- /dev/null +++ b/dpnp/backend/include/dpnp_gen_2arg_2type_tbl.hpp @@ -0,0 +1,91 @@ +//***************************************************************************** +// Copyright (c) 2023, 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. +//***************************************************************************** + +/* + * This header file contains single argument element wise functions definitions + * + * Macro `MACRO_2ARG_2TYPES_LOGIC_OP` must be defined before usage + * + * Parameters: + * - public name of the function and kernel name + * - operation used to calculate the result + * + */ + +#ifndef MACRO_2ARG_2TYPES_LOGIC_OP +#error "MACRO_2ARG_2TYPES_LOGIC_OP is not defined" +#endif + +#ifdef _SECTION_DOCUMENTATION_GENERATION_ + +#define MACRO_2ARG_2TYPES_LOGIC_OP(__name__, __operation__) \ + /** @ingroup BACKEND_API */ \ + /** @brief Per element operation function __name__ */ \ + /** */ \ + /** Function "__name__" executes operator "__operation__" over corresponding elements of input arrays */ \ + /** */ \ + /** @param[in] q_ref Reference to SYCL queue. */ \ + /** @param[out] result_out Output array. */ \ + /** @param[in] result_size Output array size. */ \ + /** @param[in] result_ndim Number of output array dimensions. */ \ + /** @param[in] result_shape Output array shape. */ \ + /** @param[in] result_strides Output array strides. */ \ + /** @param[in] input1_in Input array 1. */ \ + /** @param[in] input1_size Input array 1 size. */ \ + /** @param[in] input1_ndim Number of input array 1 dimensions. */ \ + /** @param[in] input1_shape Input array 1 shape. */ \ + /** @param[in] input1_strides Input array 1 strides. */ \ + /** @param[in] input2_in Input array 2. */ \ + /** @param[in] input2_size Input array 2 size. */ \ + /** @param[in] input2_ndim Number of input array 2 dimensions. */ \ + /** @param[in] input2_shape Input array 2 shape. */ \ + /** @param[in] input2_strides Input array 2 strides. */ \ + /** @param[in] where Where condition. */ \ + /** @param[in] dep_event_vec_ref Reference to vector of SYCL events. */ \ + template \ + DPCTLSyclEventRef __name__(DPCTLSyclQueueRef q_ref, \ + void* result_out, \ + const size_t result_size, \ + const size_t result_ndim, \ + const shape_elem_type* result_shape, \ + const shape_elem_type* result_strides, \ + const void* input1_in, \ + const size_t input1_size, \ + const size_t input1_ndim, \ + const shape_elem_type* input1_shape, \ + const shape_elem_type* input1_strides, \ + const void* input2_in, \ + const size_t input2_size, \ + const size_t input2_ndim, \ + const shape_elem_type* input2_shape, \ + const shape_elem_type* input2_strides, \ + const size_t* where, \ + const DPCTLEventVectorRef dep_event_vec_ref); + +#endif + +MACRO_2ARG_2TYPES_LOGIC_OP(dpnp_less_equal_c, input1_elem <= input2_elem) + +#undef MACRO_2ARG_2TYPES_LOGIC_OP diff --git a/dpnp/backend/include/dpnp_iface.hpp b/dpnp/backend/include/dpnp_iface.hpp index 756c7082598..713e3e82197 100644 --- a/dpnp/backend/include/dpnp_iface.hpp +++ b/dpnp/backend/include/dpnp_iface.hpp @@ -1806,6 +1806,29 @@ INP_DLLEXPORT void dpnp_invert_c(void* array1_in, void* result, size_t size); #include +#define MACRO_2ARG_2TYPES_LOGIC_OP(__name__, __operation__) \ + template \ + INP_DLLEXPORT DPCTLSyclEventRef __name__(DPCTLSyclQueueRef q_ref, \ + void* result_out, \ + const size_t result_size, \ + const size_t result_ndim, \ + const shape_elem_type* result_shape, \ + const shape_elem_type* result_strides, \ + const void* input1_in, \ + const size_t input1_size, \ + const size_t input1_ndim, \ + const shape_elem_type* input1_shape, \ + const shape_elem_type* input1_strides, \ + const void* input2_in, \ + const size_t input2_size, \ + const size_t input2_ndim, \ + const shape_elem_type* input2_shape, \ + const shape_elem_type* input2_strides, \ + const size_t* where, \ + const DPCTLEventVectorRef dep_event_vec_ref); + +#include + #define MACRO_2ARG_3TYPES_OP(__name__, __operation1__, __operation2__) \ template \ INP_DLLEXPORT DPCTLSyclEventRef __name__(DPCTLSyclQueueRef q_ref, \ diff --git a/dpnp/backend/include/dpnp_iface_fptr.hpp b/dpnp/backend/include/dpnp_iface_fptr.hpp index 126fdd24f3b..18e3629366d 100644 --- a/dpnp/backend/include/dpnp_iface_fptr.hpp +++ b/dpnp/backend/include/dpnp_iface_fptr.hpp @@ -193,6 +193,7 @@ enum class DPNPFuncName : size_t DPNP_FN_KRON_EXT, /**< Used in numpy.kron() impl, requires extra parameters */ DPNP_FN_LEFT_SHIFT, /**< Used in numpy.left_shift() impl */ DPNP_FN_LEFT_SHIFT_EXT, /**< Used in numpy.left_shift() impl, requires extra parameters */ + DPNP_FN_LESS_EQUAL_EXT, /**< Used in numpy.less_equal() impl, requires extra parameters */ DPNP_FN_LOG, /**< Used in numpy.log() impl */ DPNP_FN_LOG_EXT, /**< Used in numpy.log() impl, requires extra parameters */ DPNP_FN_LOG10, /**< Used in numpy.log10() impl */ diff --git a/dpnp/backend/kernels/dpnp_krnl_logic.cpp b/dpnp/backend/kernels/dpnp_krnl_logic.cpp index 10924691358..6be989a4ec8 100644 --- a/dpnp/backend/kernels/dpnp_krnl_logic.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_logic.cpp @@ -1,5 +1,5 @@ //***************************************************************************** -// Copyright (c) 2016-2020, Intel Corporation +// Copyright (c) 2016-2023, Intel Corporation // All rights reserved. // // Redistribution and use in source and binary forms, with or without @@ -27,6 +27,7 @@ #include "dpnp_fptr.hpp" #include "dpnp_iface.hpp" +#include "dpnp_iterator.hpp" #include "dpnpc_memory_adapter.hpp" #include "queue_sycl.hpp" @@ -286,6 +287,265 @@ DPCTLSyclEventRef (*dpnp_any_ext_c)(DPCTLSyclQueueRef, const size_t, const DPCTLEventVectorRef) = dpnp_any_c<_DataType, _ResultType>; + +#define MACRO_2ARG_2TYPES_LOGIC_OP(__name__, __operation__) \ + template \ + class __name__##_kernel; \ + \ + template \ + class __name__##_broadcast_kernel; \ + \ + template \ + class __name__##_strides_kernel; \ + \ + template \ + DPCTLSyclEventRef __name__(DPCTLSyclQueueRef q_ref, \ + void* result_out, \ + const size_t result_size, \ + const size_t result_ndim, \ + const shape_elem_type* result_shape, \ + const shape_elem_type* result_strides, \ + const void* input1_in, \ + const size_t input1_size, \ + const size_t input1_ndim, \ + const shape_elem_type* input1_shape, \ + const shape_elem_type* input1_strides, \ + const void* input2_in, \ + const size_t input2_size, \ + const size_t input2_ndim, \ + const shape_elem_type* input2_shape, \ + const shape_elem_type* input2_strides, \ + const size_t* where, \ + const DPCTLEventVectorRef dep_event_vec_ref) \ + { \ + /* avoid warning unused variable*/ \ + (void)where; \ + (void)dep_event_vec_ref; \ + \ + DPCTLSyclEventRef event_ref = nullptr; \ + \ + if (!input1_size || !input2_size) \ + { \ + return event_ref; \ + } \ + \ + sycl::queue q = *(reinterpret_cast(q_ref)); \ + \ + _DataType_input1* input1_data = static_cast<_DataType_input1 *>(const_cast(input1_in)); \ + _DataType_input2* input2_data = static_cast<_DataType_input2 *>(const_cast(input2_in)); \ + bool* result = static_cast(result_out); \ + \ + bool use_broadcasting = !array_equal(input1_shape, input1_ndim, input2_shape, input2_ndim); \ + \ + shape_elem_type* input1_shape_offsets = new shape_elem_type[input1_ndim]; \ + \ + get_shape_offsets_inkernel(input1_shape, input1_ndim, input1_shape_offsets); \ + bool use_strides = !array_equal(input1_strides, input1_ndim, input1_shape_offsets, input1_ndim); \ + delete[] input1_shape_offsets; \ + \ + shape_elem_type* input2_shape_offsets = new shape_elem_type[input2_ndim]; \ + \ + get_shape_offsets_inkernel(input2_shape, input2_ndim, input2_shape_offsets); \ + use_strides = \ + use_strides || !array_equal(input2_strides, input2_ndim, input2_shape_offsets, input2_ndim); \ + delete[] input2_shape_offsets; \ + \ + sycl::event event; \ + sycl::range<1> gws(result_size); /* used only when use_broadcasting or use_strides is true */ \ + \ + if (use_broadcasting) \ + { \ + DPNPC_id<_DataType_input1>* input1_it; \ + const size_t input1_it_size_in_bytes = sizeof(DPNPC_id<_DataType_input1>); \ + input1_it = reinterpret_cast*>(dpnp_memory_alloc_c(q_ref, \ + input1_it_size_in_bytes)); \ + new (input1_it) \ + DPNPC_id<_DataType_input1>(q_ref, input1_data, input1_shape, input1_strides, input1_ndim); \ + \ + input1_it->broadcast_to_shape(result_shape, result_ndim); \ + \ + DPNPC_id<_DataType_input2>* input2_it; \ + const size_t input2_it_size_in_bytes = sizeof(DPNPC_id<_DataType_input2>); \ + input2_it = reinterpret_cast*>(dpnp_memory_alloc_c(q_ref, \ + input2_it_size_in_bytes)); \ + new (input2_it) \ + DPNPC_id<_DataType_input2>(q_ref, input2_data, input2_shape, input2_strides, input2_ndim); \ + \ + input2_it->broadcast_to_shape(result_shape, result_ndim); \ + \ + auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { \ + const size_t i = global_id[0]; /* for (size_t i = 0; i < result_size; ++i) */ \ + { \ + const _DataType_input1 input1_elem = (*input1_it)[i]; \ + const _DataType_input2 input2_elem = (*input2_it)[i]; \ + result[i] = __operation__; \ + } \ + }; \ + auto kernel_func = [&](sycl::handler& cgh) { \ + cgh.parallel_for< \ + class __name__##_broadcast_kernel<_DataType_input1, _DataType_input2>>( \ + gws, kernel_parallel_for_func); \ + }; \ + \ + q.submit(kernel_func).wait(); \ + \ + input1_it->~DPNPC_id(); \ + input2_it->~DPNPC_id(); \ + \ + return event_ref; \ + } \ + else if (use_strides) \ + { \ + if ((result_ndim != input1_ndim) || (result_ndim != input2_ndim)) \ + { \ + throw std::runtime_error("Result ndim=" + std::to_string(result_ndim) + \ + " mismatches with either input1 ndim=" + std::to_string(input1_ndim) + \ + " or input2 ndim=" + std::to_string(input2_ndim)); \ + } \ + \ + /* memory transfer optimization, use USM-host for temporary speeds up tranfer to device */ \ + using usm_host_allocatorT = sycl::usm_allocator; \ + \ + size_t strides_size = 3 * result_ndim; \ + shape_elem_type *dev_strides_data = sycl::malloc_device(strides_size, q); \ + \ + /* create host temporary for packed strides managed by shared pointer */ \ + auto strides_host_packed = std::vector(strides_size, \ + usm_host_allocatorT(q)); \ + \ + /* packed vector is concatenation of result_strides, input1_strides and input2_strides */ \ + std::copy(result_strides, result_strides + result_ndim, strides_host_packed.begin()); \ + std::copy(input1_strides, input1_strides + result_ndim, strides_host_packed.begin() + result_ndim); \ + std::copy(input2_strides, input2_strides + result_ndim, strides_host_packed.begin() + 2 * result_ndim); \ + \ + auto copy_strides_ev = q.copy(strides_host_packed.data(), \ + dev_strides_data, \ + strides_host_packed.size()); \ + \ + auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { \ + const size_t output_id = global_id[0]; /* for (size_t i = 0; i < result_size; ++i) */ \ + { \ + const shape_elem_type *result_strides_data = &dev_strides_data[0]; \ + const shape_elem_type *input1_strides_data = &dev_strides_data[1]; \ + const shape_elem_type *input2_strides_data = &dev_strides_data[2]; \ + \ + size_t input1_id = 0; \ + size_t input2_id = 0; \ + \ + for (size_t i = 0; i < result_ndim; ++i) \ + { \ + const size_t output_xyz_id = \ + get_xyz_id_by_id_inkernel(output_id, result_strides_data, result_ndim, i); \ + input1_id += output_xyz_id * input1_strides_data[i]; \ + input2_id += output_xyz_id * input2_strides_data[i]; \ + } \ + \ + const _DataType_input1 input1_elem = input1_data[input1_id]; \ + const _DataType_input2 input2_elem = input2_data[input2_id]; \ + result[output_id] = __operation__; \ + } \ + }; \ + auto kernel_func = [&](sycl::handler& cgh) { \ + cgh.depends_on(copy_strides_ev); \ + cgh.parallel_for< \ + class __name__##_strides_kernel<_DataType_input1, _DataType_input2>>( \ + gws, kernel_parallel_for_func); \ + }; \ + \ + q.submit(kernel_func).wait(); \ + \ + sycl::free(dev_strides_data, q); \ + return event_ref; \ + } \ + else \ + { \ + constexpr size_t lws = 64; \ + constexpr unsigned int vec_sz = 8; \ + constexpr sycl::access::address_space global_space = sycl::access::address_space::global_space; \ + \ + auto gws_range = sycl::range<1>(((result_size + lws * vec_sz - 1) / (lws * vec_sz)) * lws); \ + auto lws_range = sycl::range<1>(lws); \ + \ + auto kernel_parallel_for_func = [=](sycl::nd_item<1> nd_it) { \ + auto sg = nd_it.get_sub_group(); \ + size_t start = vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + \ + sg.get_group_id()[0] * sg.get_max_local_range()[0]); \ + size_t end = start + static_cast(vec_sz); \ + \ + if (end < result_size) { \ + sycl::vec<_DataType_input1, vec_sz> x1 = \ + sg.load(sycl::multi_ptr<_DataType_input1, global_space>(&input1_data[start])); \ + sycl::vec<_DataType_input2, vec_sz> x2 = \ + sg.load(sycl::multi_ptr<_DataType_input2, global_space>(&input2_data[start])); \ + sycl::vec res_vec; \ + \ + for (size_t k = 0; k < vec_sz; ++k) { \ + const _DataType_input1 input1_elem = x1[k]; \ + const _DataType_input2 input2_elem = x2[k]; \ + res_vec[k] = __operation__; \ + } \ + sg.store(sycl::multi_ptr(&result[start]), res_vec); \ + \ + } \ + else { \ + for (size_t k = start; k < result_size; ++k) { \ + const _DataType_input1 input1_elem = input1_data[k]; \ + const _DataType_input2 input2_elem = input2_data[k]; \ + result[k] = __operation__; \ + } \ + } \ + }; \ + \ + auto kernel_func = [&](sycl::handler& cgh) { \ + cgh.parallel_for>( \ + sycl::nd_range<1>(gws_range, lws_range), kernel_parallel_for_func); \ + }; \ + event = q.submit(kernel_func); \ + } \ + \ + event_ref = reinterpret_cast(&event); \ + return DPCTLEvent_Copy(event_ref); \ + } \ + \ + template \ + DPCTLSyclEventRef (*__name__##_ext)(DPCTLSyclQueueRef, \ + void*, \ + const size_t, \ + const size_t, \ + const shape_elem_type*, \ + const shape_elem_type*, \ + const void*, \ + const size_t, \ + const size_t, \ + const shape_elem_type*, \ + const shape_elem_type*, \ + const void*, \ + const size_t, \ + const size_t, \ + const shape_elem_type*, \ + const shape_elem_type*, \ + const size_t*, \ + const DPCTLEventVectorRef) = __name__<_DataType_input1, \ + _DataType_input2>; + +#include + +template +static void func_map_logic_2arg_2type_core(func_map_t& fmap) +{ + ((fmap[DPNPFuncName::DPNP_FN_LESS_EQUAL_EXT][FT1][FTs] = + {eft_BLN, (void*)dpnp_less_equal_c_ext, func_type_map_t::find_type>}), ...); +} + +template +static void func_map_logic_2arg_2type_helper(func_map_t& fmap) +{ + ((func_map_logic_2arg_2type_core(fmap)), ...); +} + 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}; @@ -378,5 +638,7 @@ void func_map_init_logic(func_map_t& fmap) fmap[DPNPFuncName::DPNP_FN_ANY_EXT][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_any_ext_c}; fmap[DPNPFuncName::DPNP_FN_ANY_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_any_ext_c}; + func_map_logic_2arg_2type_helper(fmap); + return; } diff --git a/dpnp/backend/src/dpnp_fptr.hpp b/dpnp/backend/src/dpnp_fptr.hpp index 5b10bc71a8b..76116cafae7 100644 --- a/dpnp/backend/src/dpnp_fptr.hpp +++ b/dpnp/backend/src/dpnp_fptr.hpp @@ -1,5 +1,5 @@ //***************************************************************************** -// Copyright (c) 2016-2020, Intel Corporation +// Copyright (c) 2016-2023, Intel Corporation // All rights reserved. // // Redistribution and use in source and binary forms, with or without @@ -33,6 +33,7 @@ #define BACKEND_FPTR_H #include +#include #include @@ -64,6 +65,40 @@ const DPNPFuncType eft_C64 = DPNPFuncType::DPNP_FT_CMPLX64; const DPNPFuncType eft_C128 = DPNPFuncType::DPNP_FT_CMPLX128; const DPNPFuncType eft_BLN = DPNPFuncType::DPNP_FT_BOOL; +/** + * An internal structure to build a pair of Data type enum value with C++ type + */ +template +struct func_type_pair_t +{ + using type = T; + + static func_type_pair_t get_pair(std::integral_constant) { return {}; } +}; + +/** + * An internal structure to create a map of Data type enum value associated with C++ type + */ +template +struct func_type_map_factory_t : public Ps... +{ + using Ps::get_pair...; + + template + using find_type = typename decltype(get_pair(std::integral_constant{}))::type; +}; + +/** + * A map of the FPTR interface to link Data type enum value with accociated C++ type + */ +typedef func_type_map_factory_t, + func_type_pair_t, + func_type_pair_t, + func_type_pair_t, + func_type_pair_t, + func_type_pair_t>, + func_type_pair_t>> func_type_map_t; + /** * FPTR interface initialization functions */ diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index e0c82b6125c..0c30fa18b6f 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -1,7 +1,7 @@ # cython: language_level=3 # -*- coding: utf-8 -*- # ***************************************************************************** -# Copyright (c) 2016-2022, Intel Corporation +# Copyright (c) 2016-2023, Intel Corporation # All rights reserved. # # Redistribution and use in source and binary forms, with or without @@ -169,6 +169,7 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_KRON_EXT DPNP_FN_LEFT_SHIFT DPNP_FN_LEFT_SHIFT_EXT + DPNP_FN_LESS_EQUAL_EXT DPNP_FN_LOG DPNP_FN_LOG_EXT DPNP_FN_LOG10 @@ -429,7 +430,7 @@ ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_2in_1out_strides_t)(c_dpctl.DPCTLSyclQu const shape_elem_type * , const shape_elem_type * , const long * , - const c_dpctl.DPCTLEventVectorRef) + const c_dpctl.DPCTLEventVectorRef) except + ctypedef void(*fptr_blas_gemm_2in_1out_t)(void *, void * , void * , size_t, size_t, size_t) ctypedef c_dpctl.DPCTLSyclEventRef(*dpnp_reduction_c_t)(c_dpctl.DPCTLSyclQueueRef, void *, diff --git a/dpnp/dpnp_algo/dpnp_algo_logic.pyx b/dpnp/dpnp_algo/dpnp_algo_logic.pyx index e0b928ddf02..0aa8f949cfb 100644 --- a/dpnp/dpnp_algo/dpnp_algo_logic.pyx +++ b/dpnp/dpnp_algo/dpnp_algo_logic.pyx @@ -1,7 +1,7 @@ # cython: language_level=3 # -*- coding: utf-8 -*- # ***************************************************************************** -# Copyright (c) 2016-2020, Intel Corporation +# Copyright (c) 2016-2023, Intel Corporation # All rights reserved. # # Redistribution and use in source and binary forms, with or without @@ -286,18 +286,13 @@ cpdef utils.dpnp_descriptor dpnp_less(utils.dpnp_descriptor input1, utils.dpnp_d return result -cpdef utils.dpnp_descriptor dpnp_less_equal(utils.dpnp_descriptor input1, utils.dpnp_descriptor input2): - 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] = dpnp.bool(input1.get_pyobj()[i] <= input2.get_pyobj()[i]) +cpdef utils.dpnp_descriptor dpnp_less_equal(utils.dpnp_descriptor x1_obj, + utils.dpnp_descriptor x2_obj, + object dtype=None, + utils.dpnp_descriptor out=None, + object where=True): + return call_fptr_2in_1out_strides(DPNP_FN_LESS_EQUAL_EXT, x1_obj, x2_obj, dtype, out, where, func_name="less_equal") - return result cpdef utils.dpnp_descriptor dpnp_logical_and(utils.dpnp_descriptor input1, utils.dpnp_descriptor input2): diff --git a/dpnp/dpnp_iface_logic.py b/dpnp/dpnp_iface_logic.py index 0f1e1b5fc0e..96fa795d4d6 100644 --- a/dpnp/dpnp_iface_logic.py +++ b/dpnp/dpnp_iface_logic.py @@ -2,7 +2,7 @@ # distutils: language = c++ # -*- coding: utf-8 -*- # ***************************************************************************** -# Copyright (c) 2016-2020, Intel Corporation +# Copyright (c) 2016-2023, Intel Corporation # All rights reserved. # # Redistribution and use in source and binary forms, with or without @@ -576,17 +576,32 @@ def less(x1, x2): return call_origin(numpy.less, x1, x2) -def less_equal(x1, x2): +def less_equal(x1, + x2, + /, + out=None, + *, + where=True, + dtype=None, + subok=True): """ - Return (x1 <= x2) element-wise. + Return the truth value of (x1 <= x2) element-wise. For full documentation refer to :obj:`numpy.less_equal`. + Returns + ------- + out : dpnp.ndarray + Output array of bool type, element-wise comparison of `x1` and `x2`. + Limitations ----------- - At least either ``x1`` or ``x2`` should be as :obj:`dpnp.ndarray`. + Parameters `x1` and `x2` are supported as either :class:`dpnp.ndarray` or scalar, + but not both (at least either `x1` or `x2` should be as :class:`dpnp.ndarray`). + Parameters `out`, `where`, `dtype` and `subok` are supported with their default values. Otherwise the function will be executed sequentially on CPU. - Input array data types are limited by supported DPNP :ref:`Data types`. + Input array data types are limited by supported DPNP :ref:`Data types`, + excluding `dpnp.complex64` and `dpnp.complex128`. See Also -------- @@ -607,16 +622,25 @@ def less_equal(x1, x2): """ - # x1_desc = dpnp.get_dpnp_descriptor(x1) - # x2_desc = dpnp.get_dpnp_descriptor(x2) - # if x1_desc and x2_desc: - # if x1_desc.size < 2: - # pass - # elif x2_desc.size < 2: - # pass - # else: - # return dpnp_less_equal(x1_desc, x2_desc).get_pyobj() - + if out is not None: + pass + elif where is not True: + pass + elif dtype is not None: + pass + elif subok is not True: + pass + elif dpnp.isscalar(x1) and dpnp.isscalar(x2): + # at least either x1 or x2 has to be an array + pass + else: + # get a common queue to copy data from the host into a device if any input is scalar + queue = get_common_allocation_queue([x1, x2]) if dpnp.isscalar(x1) or dpnp.isscalar(x2) else None + + x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_strides=False, copy_when_nondefault_queue=False, alloc_queue=queue) + x2_desc = dpnp.get_dpnp_descriptor(x2, copy_when_strides=False, copy_when_nondefault_queue=False, alloc_queue=queue) + if x1_desc and x2_desc: + return dpnp_less_equal(x1_desc, x2_desc).get_pyobj() return call_origin(numpy.less_equal, x1, x2) diff --git a/dpnp/dpnp_utils/dpnp_algo_utils.pxd b/dpnp/dpnp_utils/dpnp_algo_utils.pxd index 0924dae2640..db7127319bb 100644 --- a/dpnp/dpnp_utils/dpnp_algo_utils.pxd +++ b/dpnp/dpnp_utils/dpnp_algo_utils.pxd @@ -1,7 +1,7 @@ # cython: language_level=3 # -*- coding: utf-8 -*- # ***************************************************************************** -# Copyright (c) 2016-2020, Intel Corporation +# Copyright (c) 2016-2023, Intel Corporation # All rights reserved. # # Redistribution and use in source and binary forms, with or without @@ -123,7 +123,7 @@ cdef class dpnp_descriptor: cdef void * get_data(self) -cdef shape_type_c get_common_shape(shape_type_c input1_shape, shape_type_c input2_shape) +cdef shape_type_c get_common_shape(shape_type_c input1_shape, shape_type_c input2_shape) except * """ Calculate common shape from input shapes """ diff --git a/dpnp/dpnp_utils/dpnp_algo_utils.pyx b/dpnp/dpnp_utils/dpnp_algo_utils.pyx index 2e04dd96bd6..6605770be62 100644 --- a/dpnp/dpnp_utils/dpnp_algo_utils.pyx +++ b/dpnp/dpnp_utils/dpnp_algo_utils.pyx @@ -429,7 +429,9 @@ cpdef find_common_type(object x1_obj, object x2_obj): return numpy.find_common_type(array_types, scalar_types) -cdef shape_type_c get_common_shape(shape_type_c input1_shape, shape_type_c input2_shape): +cdef shape_type_c get_common_shape(shape_type_c input1_shape, shape_type_c input2_shape) except *: + cdef shape_type_c input1_shape_orig = input1_shape + cdef shape_type_c input2_shape_orig = input2_shape cdef shape_type_c result_shape # ex (8, 1, 6, 1) and (7, 1, 5) -> (8, 1, 6, 1) and (1, 7, 1, 5) @@ -446,9 +448,9 @@ cdef shape_type_c get_common_shape(shape_type_c input1_shape, shape_type_c input elif input2_shape[it] == 1: result_shape.push_back(input1_shape[it]) else: - err_msg = f"{ERROR_PREFIX} in function get_common_shape()" - err_msg += f"operands could not be broadcast together with shapes {input1_shape} {input2_shape}" - ValueError(err_msg) + err_msg = f"{ERROR_PREFIX} in function get_common_shape(): " + err_msg += f"operands could not be broadcast together with shapes {input1_shape_orig} {input2_shape_orig}" + raise ValueError(err_msg) return result_shape diff --git a/tests/test_logic.py b/tests/test_logic.py index 7fefe91826f..d79ffaa744f 100644 --- a/tests/test_logic.py +++ b/tests/test_logic.py @@ -3,6 +3,10 @@ import dpnp import numpy +from numpy.testing import ( + assert_allclose, + assert_equal +) @pytest.mark.parametrize("type", @@ -31,11 +35,11 @@ def test_all(type, shape): np_res = numpy.all(a) dpnp_res = dpnp.all(ia) - numpy.testing.assert_allclose(dpnp_res, np_res) + assert_allclose(dpnp_res, np_res) np_res = a.all() dpnp_res = ia.all() - numpy.testing.assert_allclose(dpnp_res, np_res) + assert_allclose(dpnp_res, np_res) @pytest.mark.parametrize("type", @@ -51,7 +55,7 @@ def test_allclose(type): np_res = numpy.allclose(a, b) dpnp_res = dpnp.allclose(dpnp_a, dpnp_b) - numpy.testing.assert_allclose(dpnp_res, np_res) + assert_allclose(dpnp_res, np_res) a[0] = numpy.inf @@ -59,7 +63,7 @@ def test_allclose(type): np_res = numpy.allclose(a, b) dpnp_res = dpnp.allclose(dpnp_a, dpnp_b) - numpy.testing.assert_allclose(dpnp_res, np_res) + assert_allclose(dpnp_res, np_res) @pytest.mark.parametrize("type", @@ -88,11 +92,11 @@ def test_any(type, shape): np_res = numpy.any(a) dpnp_res = dpnp.any(ia) - numpy.testing.assert_allclose(dpnp_res, np_res) + assert_allclose(dpnp_res, np_res) np_res = a.any() dpnp_res = ia.any() - numpy.testing.assert_allclose(dpnp_res, np_res) + assert_allclose(dpnp_res, np_res) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @@ -102,7 +106,7 @@ def test_greater(): for i in range(len(a) + 1): np_res = (a > i) dpnp_res = (ia > i) - numpy.testing.assert_equal(dpnp_res, np_res) + assert_equal(dpnp_res, np_res) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @@ -112,7 +116,7 @@ def test_greater_equal(): for i in range(len(a) + 1): np_res = (a >= i) dpnp_res = (ia >= i) - numpy.testing.assert_equal(dpnp_res, np_res) + assert_equal(dpnp_res, np_res) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @@ -122,17 +126,16 @@ def test_less(): for i in range(len(a) + 1): np_res = (a < i) dpnp_res = (ia < i) - numpy.testing.assert_equal(dpnp_res, np_res) + assert_equal(dpnp_res, np_res) -@pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_less_equal(): a = numpy.array([1, 2, 3, 4, 5, 6, 7, 8]) ia = dpnp.array(a) for i in range(len(a) + 1): np_res = (a <= i) dpnp_res = (ia <= i) - numpy.testing.assert_equal(dpnp_res, np_res) + assert_equal(dpnp_res, np_res) @pytest.mark.usefixtures("allow_fall_back_on_numpy") @@ -142,4 +145,51 @@ def test_not_equal(): for i in range(len(a)): np_res = (a != i) dpnp_res = (ia != i) - numpy.testing.assert_equal(dpnp_res, np_res) + assert_equal(dpnp_res, np_res) + +@pytest.mark.parametrize("op", + ['less_equal'], + ids=['less_equal']) +@pytest.mark.parametrize("x1", + [[3, 4, 5, 6], [[1, 2, 3, 4], [5, 6, 7, 8]], [[1, 2, 5, 6], [3, 4, 7, 8], [1, 2, 7, 8]]], + ids=['[3, 4, 5, 6]', '[[1, 2, 3, 4], [5, 6, 7, 8]]', '[[1, 2, 5, 6], [3, 4, 7, 8], [1, 2, 7, 8]]']) +@pytest.mark.parametrize("x2", + [5, [1, 2, 5, 6]], + ids=['5', '[1, 2, 5, 6]']) +def test_elemwise_comparison(op, x1, x2): + create_func = lambda xp, a: xp.asarray(a) if not numpy.isscalar(a) else a + + np_x1, np_x2 = create_func(numpy, x1), create_func(numpy, x2) + dp_x1, dp_x2 = create_func(dpnp, np_x1), create_func(dpnp, np_x2) + + # x1 OP x2 + np_res = getattr(numpy, op)(np_x1, np_x2) + dpnp_res = getattr(dpnp, op)(dp_x1, dp_x2) + assert_equal(dpnp_res, np_res) + + # x2 OP x1 + np_res = getattr(numpy, op)(np_x2, np_x1) + dpnp_res = getattr(dpnp, op)(dp_x2, dp_x1) + assert_equal(dpnp_res, np_res) + + # x1[::-1] OP x2 + np_res = getattr(numpy, op)(np_x1[::-1], np_x2) + dpnp_res = getattr(dpnp, op)(dp_x1[::-1], dp_x2) + assert_equal(dpnp_res, np_res) + +@pytest.mark.parametrize("op", + ['less_equal'], + ids=['less_equal']) +@pytest.mark.parametrize("sh1", + [[10], [8, 4], [4, 1, 2]], + ids=['(10,)', '(8, 4)', '(4, 1, 2)']) +@pytest.mark.parametrize("sh2", + [[12], [4, 8], [1, 8, 6]], + ids=['(12,)', '(4, 8)', '(1, 8, 6)']) +def test_comparison_no_broadcast_with_shapes(op, sh1, sh2): + x1, x2 = dpnp.random.randn(*sh1), dpnp.random.randn(*sh2) + + # x1 OP x2 + with pytest.raises(ValueError): + getattr(dpnp, op)(x1, x2) + getattr(numpy, op)(x1.asnumpy(), x2.asnumpy()) diff --git a/tests/third_party/cupy/logic_tests/test_comparison.py b/tests/third_party/cupy/logic_tests/test_comparison.py index 0be9eaeee61..4afcea568ff 100644 --- a/tests/third_party/cupy/logic_tests/test_comparison.py +++ b/tests/third_party/cupy/logic_tests/test_comparison.py @@ -8,10 +8,10 @@ from tests.third_party.cupy import testing -@pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.gpu class TestComparison(unittest.TestCase): + @pytest.mark.usefixtures("allow_fall_back_on_numpy") @testing.for_all_dtypes(no_complex=True) @testing.numpy_cupy_allclose(atol=1e-5) def check_binary(self, name, xp, dtype): @@ -19,21 +19,26 @@ def check_binary(self, name, xp, dtype): b = testing.shaped_reverse_arange((2, 3), xp, dtype) return getattr(xp, name)(a, b) + @pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_greater(self): self.check_binary('greater') + @pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_greater_equal(self): self.check_binary('greater_equal') + @pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_less(self): self.check_binary('less') def test_less_equal(self): self.check_binary('less_equal') + @pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_not_equal(self): self.check_binary('not_equal') + @pytest.mark.usefixtures("allow_fall_back_on_numpy") def test_equal(self): self.check_binary('equal') diff --git a/tests_external/skipped_tests_numpy.tbl b/tests_external/skipped_tests_numpy.tbl index 30b66da5e66..c2c0dc78ec5 100644 --- a/tests_external/skipped_tests_numpy.tbl +++ b/tests_external/skipped_tests_numpy.tbl @@ -318,83 +318,6 @@ tests/test_datetime.py::TestDateTime::test_timedelta_np_int_construction[Y] tests/test_datetime.py::TestDateTime::test_timedelta_object_array_conversion tests/test_datetime.py::TestDateTime::test_timedelta_scalar_construction tests/test_datetime.py::TestDateTime::test_timedelta_scalar_construction_units -tests/test_defchararray.py::TestBasic::test_from_object_array -tests/test_defchararray.py::TestBasic::test_from_object_array_unicode -tests/test_defchararray.py::TestBasic::test_from_string -tests/test_defchararray.py::TestBasic::test_from_string_array -tests/test_defchararray.py::TestBasic::test_from_unicode -tests/test_defchararray.py::TestBasic::test_from_unicode_array -tests/test_defchararray.py::TestBasic::test_unicode_upconvert -tests/test_defchararray.py::TestChar::test_it -tests/test_defchararray.py::TestComparisonsMixed1::test_equal -tests/test_defchararray.py::TestComparisonsMixed1::test_greater -tests/test_defchararray.py::TestComparisonsMixed1::test_greater_equal -tests/test_defchararray.py::TestComparisonsMixed1::test_less -tests/test_defchararray.py::TestComparisonsMixed1::test_less_equal -tests/test_defchararray.py::TestComparisonsMixed1::test_not_equal -tests/test_defchararray.py::TestComparisonsMixed2::test_equal -tests/test_defchararray.py::TestComparisonsMixed2::test_greater -tests/test_defchararray.py::TestComparisonsMixed2::test_greater_equal -tests/test_defchararray.py::TestComparisonsMixed2::test_less -tests/test_defchararray.py::TestComparisonsMixed2::test_less_equal -tests/test_defchararray.py::TestComparisonsMixed2::test_not_equal -tests/test_defchararray.py::TestComparisons::test_equal -tests/test_defchararray.py::TestComparisons::test_greater -tests/test_defchararray.py::TestComparisons::test_greater_equal -tests/test_defchararray.py::TestComparisons::test_less -tests/test_defchararray.py::TestComparisons::test_less_equal -tests/test_defchararray.py::TestComparisons::test_not_equal -tests/test_defchararray.py::test_empty_indexing -tests/test_defchararray.py::TestInformation::test_count -tests/test_defchararray.py::TestInformation::test_endswith -tests/test_defchararray.py::TestInformation::test_find -tests/test_defchararray.py::TestInformation::test_index -tests/test_defchararray.py::TestInformation::test_isalnum -tests/test_defchararray.py::TestInformation::test_isalpha -tests/test_defchararray.py::TestInformation::test_isdigit -tests/test_defchararray.py::TestInformation::test_islower -tests/test_defchararray.py::TestInformation::test_isspace -tests/test_defchararray.py::TestInformation::test_istitle -tests/test_defchararray.py::TestInformation::test_isupper -tests/test_defchararray.py::TestInformation::test_len -tests/test_defchararray.py::TestInformation::test_rfind -tests/test_defchararray.py::TestInformation::test_rindex -tests/test_defchararray.py::TestInformation::test_startswith -tests/test_defchararray.py::TestMethods::test_capitalize -tests/test_defchararray.py::TestMethods::test_center -tests/test_defchararray.py::TestMethods::test_decode -tests/test_defchararray.py::TestMethods::test_encode -tests/test_defchararray.py::TestMethods::test_expandtabs -tests/test_defchararray.py::TestMethods::test_isdecimal -tests/test_defchararray.py::TestMethods::test_isnumeric -tests/test_defchararray.py::TestMethods::test_join -tests/test_defchararray.py::TestMethods::test_ljust -tests/test_defchararray.py::TestMethods::test_lower -tests/test_defchararray.py::TestMethods::test_lstrip -tests/test_defchararray.py::TestMethods::test_partition -tests/test_defchararray.py::TestMethods::test_replace -tests/test_defchararray.py::TestMethods::test_rjust -tests/test_defchararray.py::TestMethods::test_rpartition -tests/test_defchararray.py::TestMethods::test_rsplit -tests/test_defchararray.py::TestMethods::test_rstrip -tests/test_defchararray.py::TestMethods::test_split -tests/test_defchararray.py::TestMethods::test_splitlines -tests/test_defchararray.py::TestMethods::test_strip -tests/test_defchararray.py::TestMethods::test_swapcase -tests/test_defchararray.py::TestMethods::test_title -tests/test_defchararray.py::TestMethods::test_upper -tests/test_defchararray.py::TestOperations::test_add -tests/test_defchararray.py::TestOperations::test_mod -tests/test_defchararray.py::TestOperations::test_mul -tests/test_defchararray.py::TestOperations::test_radd -tests/test_defchararray.py::TestOperations::test_rmod -tests/test_defchararray.py::TestOperations::test_rmul -tests/test_defchararray.py::TestOperations::test_slice -tests/test_defchararray.py::TestVecString::test_invalid_args_tuple -tests/test_defchararray.py::TestVecString::test_invalid_function_args -tests/test_defchararray.py::TestVecString::test_invalid_result_type -tests/test_defchararray.py::TestVecString::test_non_string_array -tests/test_defchararray.py::TestWhitespace::test1 tests/test_deprecations.py::TestAlen::test_alen tests/test_deprecations.py::TestArrayDataAttributeAssignmentDeprecation::test_data_attr_assignment tests/test_deprecations.py::TestBinaryReprInsufficientWidthParameterForRepresentation::test_insufficient_width_negative