From 2ac7f3921b7838df994e6b483ef9937dcf31edc6 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Thu, 24 Aug 2023 17:51:33 +0200 Subject: [PATCH 1/6] Added support of dpnp.allclose() for a device without fp64 aspect --- dpnp/backend/kernels/dpnp_krnl_logic.cpp | 116 +++++++++++++----- dpnp/dpnp_iface_logic.py | 68 +++++++--- tests/skipped_tests.tbl | 6 +- tests/skipped_tests_gpu.tbl | 6 +- tests/test_logic.py | 3 - .../cupy/logic_tests/test_comparison.py | 19 +-- 6 files changed, 144 insertions(+), 74 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_logic.cpp b/dpnp/backend/kernels/dpnp_krnl_logic.cpp index 1757d053416a..0c49c4bf2667 100644 --- a/dpnp/backend/kernels/dpnp_krnl_logic.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_logic.cpp @@ -74,7 +74,7 @@ DPCTLSyclEventRef dpnp_all_c(DPCTLSyclQueueRef q_ref, 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_group(); + 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) + @@ -127,8 +127,72 @@ DPCTLSyclEventRef (*dpnp_all_ext_c)(DPCTLSyclQueueRef, const DPCTLEventVectorRef) = dpnp_all_c<_DataType, _ResultType>; -template -class dpnp_allclose_c_kernel; +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; + } + } + partial &= (std::abs(array1[i] - array2[i]) <= + (atol_val + rtol_val * std::abs(array2[i]))); + } + 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, @@ -140,6 +204,9 @@ DPCTLSyclEventRef dpnp_allclose_c(DPCTLSyclQueueRef q_ref, 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; @@ -152,40 +219,21 @@ DPCTLSyclEventRef dpnp_allclose_c(DPCTLSyclQueueRef q_ref, sycl::queue q = *(reinterpret_cast(q_ref)); sycl::event event; - DPNPC_ptr_adapter<_DataType1> input1_ptr(q_ref, array1_in, size); - DPNPC_ptr_adapter<_DataType2> input2_ptr(q_ref, array2_in, size); - DPNPC_ptr_adapter<_ResultType> result1_ptr(q_ref, result1, 1, true, true); - const _DataType1 *array1 = input1_ptr.get_ptr(); - const _DataType2 *array2 = input2_ptr.get_ptr(); - _ResultType *result = result1_ptr.get_ptr(); - - result[0] = true; + const _DataType1 *array1 = static_cast(array1_in); + const _DataType2 *array2 = static_cast(array2_in); + bool *result = static_cast(result1); - if (!size) { - return event_ref; + 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); } - - sycl::range<1> gws(size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - size_t i = global_id[0]; - - if (std::abs(array1[i] - array2[i]) > - (atol_val + rtol_val * std::abs(array2[i]))) - { - result[0] = false; - } - }; - - auto kernel_func = [&](sycl::handler &cgh) { - cgh.parallel_for< - class dpnp_allclose_c_kernel<_DataType1, _DataType2, _ResultType>>( - gws, kernel_parallel_for_func); - }; - - event = q.submit(kernel_func); event_ref = reinterpret_cast(&event); - return DPCTLEvent_Copy(event_ref); } @@ -269,7 +317,7 @@ DPCTLSyclEventRef dpnp_any_c(DPCTLSyclQueueRef q_ref, 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_group(); + 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) + diff --git a/dpnp/dpnp_iface_logic.py b/dpnp/dpnp_iface_logic.py index d867bbbd6609..20a5b28bd630 100644 --- a/dpnp/dpnp_iface_logic.py +++ b/dpnp/dpnp_iface_logic.py @@ -152,42 +152,74 @@ def all(x, /, axis=None, out=None, keepdims=False, *, where=True): ) -def allclose(x1, x2, rtol=1.0e-5, atol=1.0e-8, **kwargs): +def allclose(a, b, rtol=1.0e-5, atol=1.0e-8, **kwargs): """ Returns True if two arrays are element-wise equal within a tolerance. For full documentation refer to :obj:`numpy.allclose`. + Returns + ------- + out : dpnp.ndarray + A boolean 0-dim array. If its value is ``True``, + two arrays are element-wise equal within a tolerance. + Limitations ----------- - Parameters `x1` and `x2` are supported as either :class:`dpnp.ndarray` or scalar. + 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. - Input array data types are limited by supported DPNP :ref:`Data types`. + Parameters `rtol` and `atol` are supported as scalars. Otherwise + ``TypeError`` exeption will be raised. + Input array data types are limited by supported integer and + floating DPNP :ref:`Data types`. + + See Also + -------- + :obj:`dpnp.isclose` : Test whether two arrays are element-wise equal. + :obj:`dpnp.all` : Test whether all elements evaluate to True. + :obj:`dpnp.any` : Test whether any element evaluates to True. + :obj:`dpnp.equal` : Return (x1 == x2) element-wise. Examples -------- >>> import dpnp as np - >>> np.allclose([1e10,1e-7], [1.00001e10,1e-8]) - >>> False + >>> np.allclose(np.array([1e10, 1e-7]), np.array([1.00001e10, 1e-8])) + array([False]) + >>> np.allclose(np.array([1.0, np.nan]), np.array([1.0, np.nan])) + array([False]) + >>> np.allclose(np.array([1.0, np.inf]), np.array([1.0, np.inf])) + array([ True]) """ - rtol_is_scalar = dpnp.isscalar(rtol) - atol_is_scalar = dpnp.isscalar(atol) - x1_desc = dpnp.get_dpnp_descriptor(x1, copy_when_nondefault_queue=False) - x2_desc = dpnp.get_dpnp_descriptor(x2, copy_when_nondefault_queue=False) - - if x1_desc and x2_desc and not kwargs: - if not rtol_is_scalar or not atol_is_scalar: - pass - else: - result_obj = dpnp_allclose(x1_desc, x2_desc, rtol, atol).get_pyobj() - result = dpnp.convert_single_elem_array_to_scalar(result_obj) + if dpnp.isscalar(a) and dpnp.isscalar(b): + # at least one of inputs has to be an array + pass + elif kwargs: + pass + else: + if not dpnp.isscalar(rtol): + raise TypeError( + "An argument `rtol` must be a scalar, but got {}".format( + type(rtol) + ) + ) + elif not dpnp.isscalar(atol): + raise TypeError( + "An argument `atol` must be a scalar, but got {}".format( + type(atol) + ) + ) - return result + 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, x1, x2, rtol=rtol, atol=atol, **kwargs) + return call_origin(numpy.allclose, a, b, rtol=rtol, atol=atol, **kwargs) def any(x, /, axis=None, out=None, keepdims=False, *, where=True): diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl index c7627f9c9d0c..cc255790e98d 100644 --- a/tests/skipped_tests.tbl +++ b/tests/skipped_tests.tbl @@ -452,11 +452,7 @@ tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_transpose tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_transposed_tensordot_with_int_axes tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_transposed_tensordot_with_list_axes tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_reversed_vdot -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_array_scalar -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_finite -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_infinite -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_infinite_equal_nan -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_min_int + tests/third_party/cupy/logic_tests/test_comparison.py::TestArrayEqual::test_array_equal_broadcast_not_allowed tests/third_party/cupy/logic_tests/test_comparison.py::TestArrayEqual::test_array_equal_diff_dtypes_is_equal tests/third_party/cupy/logic_tests/test_comparison.py::TestArrayEqual::test_array_equal_diff_dtypes_not_equal diff --git a/tests/skipped_tests_gpu.tbl b/tests/skipped_tests_gpu.tbl index 5bada1df5c8a..663c3f1d90e5 100644 --- a/tests/skipped_tests_gpu.tbl +++ b/tests/skipped_tests_gpu.tbl @@ -598,11 +598,7 @@ tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_transpose tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_tensordot_zero_dim tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_transposed_dot_with_out_f_contiguous tests/third_party/cupy/linalg_tests/test_product.py::TestProduct::test_transposed_multidim_vdot -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_array_scalar -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_finite -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_infinite -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_infinite_equal_nan -tests/third_party/cupy/logic_tests/test_comparison.py::TestAllclose::test_allclose_min_int + tests/third_party/cupy/logic_tests/test_comparison.py::TestArrayEqual::test_array_equal_broadcast_not_allowed tests/third_party/cupy/logic_tests/test_comparison.py::TestArrayEqual::test_array_equal_diff_dtypes_is_equal tests/third_party/cupy/logic_tests/test_comparison.py::TestArrayEqual::test_array_equal_diff_dtypes_not_equal diff --git a/tests/test_logic.py b/tests/test_logic.py index 7be9e6e1ac8c..7eacd0bc2595 100644 --- a/tests/test_logic.py +++ b/tests/test_logic.py @@ -44,9 +44,6 @@ def test_all(type, shape): assert_allclose(dpnp_res, np_res) -@pytest.mark.skipif( - not has_support_aspect64(), reason="Aborted on Iris Xe: SAT-5988" -) @pytest.mark.parametrize("type", get_all_dtypes(no_bool=True, no_complex=True)) def test_allclose(type): a = numpy.random.rand(10) diff --git a/tests/third_party/cupy/logic_tests/test_comparison.py b/tests/third_party/cupy/logic_tests/test_comparison.py index a37479c4fe21..739ff9f7c959 100644 --- a/tests/third_party/cupy/logic_tests/test_comparison.py +++ b/tests/third_party/cupy/logic_tests/test_comparison.py @@ -121,15 +121,15 @@ class TestAllclose(unittest.TestCase): @testing.for_all_dtypes() @testing.numpy_cupy_equal() def test_allclose_finite(self, xp, dtype): - 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.allclose(a, b) @testing.for_all_dtypes() @testing.numpy_cupy_equal() def test_allclose_min_int(self, xp, dtype): - 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.allclose(a, b) @testing.for_float_dtypes() @@ -138,24 +138,25 @@ def test_allclose_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.allclose(a, b) @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") 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.allclose(a, b, equal_nan=True) @testing.for_all_dtypes() @testing.numpy_cupy_equal() def test_allclose_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.allclose(a, b) From ec0f82318827a0beada55fb1fe49fce2456ad8db Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Fri, 25 Aug 2023 15:51:54 +0200 Subject: [PATCH 2/6] Added tests for SYCL queue and USM type --- .github/workflows/conda-package.yml | 1 + dpnp/dpnp_iface_logic.py | 21 ++++++++++++++++++--- tests/skipped_tests_gpu_no_fp64.tbl | 4 ++-- tests/test_sycl_queue.py | 5 +++++ tests/test_usm_type.py | 5 +++++ 5 files changed, 31 insertions(+), 5 deletions(-) diff --git a/.github/workflows/conda-package.yml b/.github/workflows/conda-package.yml index b522fd272895..ebb1697ae378 100644 --- a/.github/workflows/conda-package.yml +++ b/.github/workflows/conda-package.yml @@ -26,6 +26,7 @@ env: test_umath.py test_usm_type.py third_party/cupy/linalg_tests/test_product.py + third_party/cupy/logic_tests/test_comparison.py third_party/cupy/logic_tests/test_truth.py third_party/cupy/manipulation_tests/test_basic.py third_party/cupy/manipulation_tests/test_join.py diff --git a/dpnp/dpnp_iface_logic.py b/dpnp/dpnp_iface_logic.py index 20a5b28bd630..7f8b42c6a25d 100644 --- a/dpnp/dpnp_iface_logic.py +++ b/dpnp/dpnp_iface_logic.py @@ -186,11 +186,19 @@ def allclose(a, b, rtol=1.0e-5, atol=1.0e-8, **kwargs): Examples -------- >>> import dpnp as np - >>> np.allclose(np.array([1e10, 1e-7]), np.array([1.00001e10, 1e-8])) + >>> a = np.array([1e10, 1e-7]) + >>> b = np.array([1.00001e10, 1e-8]) + >>> np.allclose(a, b) array([False]) - >>> np.allclose(np.array([1.0, np.nan]), np.array([1.0, np.nan])) + + >>> a = np.array([1.0, np.nan]) + >>> b = np.array([1.0, np.nan]) + >>> np.allclose(a, b) array([False]) - >>> np.allclose(np.array([1.0, np.inf]), np.array([1.0, np.inf])) + + >>> a = np.array([1.0, np.inf]) + >>> b = np.array([1.0, np.inf]) + >>> np.allclose(a, b) array([ True]) """ @@ -214,6 +222,13 @@ def allclose(a, b, rtol=1.0e-5, atol=1.0e-8, **kwargs): ) ) + if dpnp.isscalar(a): + a = dpnp.full_like(b, fill_value=a) + elif dpnp.isscalar(b): + b = dpnp.full_like(a, fill_value=b) + elif a.shape != b.shape: + a, b = dpt.broadcast_arrays(a, b) + 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: diff --git a/tests/skipped_tests_gpu_no_fp64.tbl b/tests/skipped_tests_gpu_no_fp64.tbl index 31c4b499cd0a..e09257922d1b 100644 --- a/tests/skipped_tests_gpu_no_fp64.tbl +++ b/tests/skipped_tests_gpu_no_fp64.tbl @@ -443,8 +443,8 @@ tests/test_sycl_queue.py::test_array_creation[opencl:gpu:0-arange-arg0-kwargs0] tests/test_sycl_queue.py::test_array_creation[level_zero:gpu:0-arange-arg0-kwargs0] tests/test_sycl_queue.py::test_1in_1out[opencl:gpu:0-gradient-data10] tests/test_sycl_queue.py::test_1in_1out[level_zero:gpu:0-gradient-data10] -tests/test_sycl_queue.py::test_2in_1out[opencl:gpu:0-power-data112-data212] -tests/test_sycl_queue.py::test_2in_1out[level_zero:gpu:0-power-data112-data212] +tests/test_sycl_queue.py::test_2in_1out[opencl:gpu:0-power-data113-data213] +tests/test_sycl_queue.py::test_2in_1out[level_zero:gpu:0-power-data113-data213] tests/test_sycl_queue.py::test_out_2in_1out[opencl:gpu:0-power-data19-data29] tests/test_sycl_queue.py::test_out_2in_1out[level_zero:gpu:0-power-data19-data29] tests/test_sycl_queue.py::test_eig[opencl:gpu:0] diff --git a/tests/test_sycl_queue.py b/tests/test_sycl_queue.py index 78f98c61581c..71ee6e532101 100644 --- a/tests/test_sycl_queue.py +++ b/tests/test_sycl_queue.py @@ -276,6 +276,11 @@ def test_1in_1out(func, data, device): [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0], [0.0, 1.0, 2.0, 0.0, 1.0, 2.0, 0.0, 1.0, 2.0], ), + pytest.param( + "allclose", + [1.0, dpnp.inf, -dpnp.inf], + [1.0, dpnp.inf, -dpnp.inf], + ), pytest.param("copysign", [0.0, 1.0, 2.0], [-1.0, 0.0, 1.0]), pytest.param("cross", [1.0, 2.0, 3.0], [4.0, 5.0, 6.0]), pytest.param( diff --git a/tests/test_usm_type.py b/tests/test_usm_type.py index aaf9775954ce..c77a8e05c381 100644 --- a/tests/test_usm_type.py +++ b/tests/test_usm_type.py @@ -295,6 +295,11 @@ def test_1in_1out(func, data, usm_type): @pytest.mark.parametrize( "func,data1,data2", [ + pytest.param( + "allclose", + [[1.2, -0.0], [-7, 2.34567]], + [[1.2, 0.0], [-7, 2.34567]], + ), pytest.param( "dot", [[0.0, 1.0, 2.0], [3.0, 4.0, 5.0]], From ea464f5735530aaee4a1d79d9b1b107468ada0e9 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Fri, 25 Aug 2023 18:36:55 +0200 Subject: [PATCH 3/6] Handled a corner case with abs(MIN_INT) --- dpnp/backend/kernels/dpnp_krnl_logic.cpp | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_logic.cpp b/dpnp/backend/kernels/dpnp_krnl_logic.cpp index 0c49c4bf2667..499a252a35d0 100644 --- a/dpnp/backend/kernels/dpnp_krnl_logic.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_logic.cpp @@ -174,8 +174,15 @@ static sycl::event dpnp_allclose(sycl::queue &q, continue; } } - partial &= (std::abs(array1[i] - array2[i]) <= - (atol_val + rtol_val * std::abs(array2[i]))); + + // casting integeral 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); From c98d4bffc442d6c852816875a8f6cc0dca165f09 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Sat, 26 Aug 2023 14:49:25 +0200 Subject: [PATCH 4/6] Increased test coverage --- dpnp/dpnp_iface.py | 23 +++++++++++ dpnp/dpnp_iface_linearalgebra.py | 5 ++- dpnp/dpnp_iface_logic.py | 7 +++- tests/test_logic.py | 65 ++++++++++++++++++++++++++++++-- 4 files changed, 93 insertions(+), 7 deletions(-) diff --git a/dpnp/dpnp_iface.py b/dpnp/dpnp_iface.py index 066b28e8e2cd..5060ace272c4 100644 --- a/dpnp/dpnp_iface.py +++ b/dpnp/dpnp_iface.py @@ -66,6 +66,7 @@ "get_normalized_queue_device", "get_usm_ndarray", "get_usm_ndarray_or_scalar", + "is_supported_array_or_scalar", "is_supported_array_type", ] @@ -453,6 +454,28 @@ def get_usm_ndarray_or_scalar(a): return a if isscalar(a) else get_usm_ndarray(a) +def is_supported_array_or_scalar(a): + """ + Return ``True`` if `a` is a scalar or an array of either + :class:`dpnp.ndarray` or :class:`dpctl.tensor.usm_ndarray` type, + ``False`` otherwise. + + Parameters + ---------- + a : array + A input scalar or an array to check the type of. + + Returns + ------- + out : bool + ``True`` if input `a` is a scalar or an array of supportedtype, + ``False`` otherwise. + + """ + + return isscalar(a) or is_supported_array_type(a) + + def is_supported_array_type(a): """ Return ``True`` if an array of either type :class:`dpnp.ndarray` diff --git a/dpnp/dpnp_iface_linearalgebra.py b/dpnp/dpnp_iface_linearalgebra.py index 70e2a8048c0e..3e3e92d82450 100644 --- a/dpnp/dpnp_iface_linearalgebra.py +++ b/dpnp/dpnp_iface_linearalgebra.py @@ -358,14 +358,15 @@ def outer(x1, x2, out=None): [1, 2, 3]]) """ + x1_is_scalar = dpnp.isscalar(x1) x2_is_scalar = dpnp.isscalar(x2) if x1_is_scalar and x2_is_scalar: pass - elif not (x1_is_scalar or dpnp.is_supported_array_type(x1)): + elif not dpnp.is_supported_array_or_scalar(x1): pass - elif not (x2_is_scalar or dpnp.is_supported_array_type(x2)): + elif not dpnp.is_supported_array_or_scalar(x2): pass else: x1_in = ( diff --git a/dpnp/dpnp_iface_logic.py b/dpnp/dpnp_iface_logic.py index 2f755ef3a213..c4704d09f9a0 100644 --- a/dpnp/dpnp_iface_logic.py +++ b/dpnp/dpnp_iface_logic.py @@ -206,6 +206,11 @@ def allclose(a, b, rtol=1.0e-5, atol=1.0e-8, **kwargs): if dpnp.isscalar(a) and dpnp.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: @@ -227,7 +232,7 @@ def allclose(a, b, rtol=1.0e-5, atol=1.0e-8, **kwargs): elif dpnp.isscalar(b): b = dpnp.full_like(a, fill_value=b) elif a.shape != b.shape: - a, b = dpt.broadcast_arrays(a, b) + 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) diff --git a/tests/test_logic.py b/tests/test_logic.py index 7eacd0bc2595..b9d2a9b4303a 100644 --- a/tests/test_logic.py +++ b/tests/test_logic.py @@ -44,13 +44,13 @@ def test_all(type, shape): assert_allclose(dpnp_res, np_res) -@pytest.mark.parametrize("type", get_all_dtypes(no_bool=True, no_complex=True)) -def test_allclose(type): +@pytest.mark.parametrize("dtype", get_all_dtypes(no_bool=True, no_complex=True)) +def test_allclose(dtype): a = numpy.random.rand(10) b = a + numpy.random.rand(10) * 1e-8 - dpnp_a = dpnp.array(a, dtype=type) - dpnp_b = dpnp.array(b, dtype=type) + dpnp_a = dpnp.array(a, dtype=dtype) + dpnp_b = dpnp.array(b, dtype=dtype) np_res = numpy.allclose(a, b) dpnp_res = dpnp.allclose(dpnp_a, dpnp_b) @@ -65,6 +65,63 @@ def test_allclose(type): assert_allclose(dpnp_res, np_res) +class TestAllClose: + @pytest.mark.parametrize("val", [1.0, 3, numpy.inf, -numpy.inf, numpy.nan]) + def test_input_0d(self, val): + dp_arr = dpnp.array(val) + np_arr = numpy.array(val) + + # array & scalar + dp_res = dpnp.allclose(dp_arr, val) + np_res = numpy.allclose(np_arr, val) + assert_allclose(dp_res, np_res) + + # scalar & array + dp_res = dpnp.allclose(val, dp_arr) + np_res = numpy.allclose(val, np_arr) + assert_allclose(dp_res, np_res) + + # two arrays + dp_res = dpnp.allclose(dp_arr, dp_arr) + np_res = numpy.allclose(np_arr, np_arr) + assert_allclose(dp_res, np_res) + + @pytest.mark.parametrize("sh_a", [(10,), (10, 10)]) + @pytest.mark.parametrize("sh_b", [(1, 10), (1, 10, 1)]) + def test_broadcast(self, sh_a, sh_b): + dp_a = dpnp.ones(sh_a) + dp_b = dpnp.ones(sh_b) + + np_a = numpy.ones(sh_a) + np_b = numpy.ones(sh_b) + + dp_res = dpnp.allclose(dp_a, dp_b) + np_res = numpy.allclose(np_a, np_b) + assert_allclose(dp_res, np_res) + + def test_input_as_scalars(self): + with pytest.raises(NotImplementedError): + 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): + dpnp.allclose(val, val) + + @pytest.mark.parametrize( + "tol", [[0.001], (1.0e-6,), dpnp.array(1.0e-3), 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" + ): + dpnp.allclose(a, b, **kw) + + @pytest.mark.parametrize("type", get_all_dtypes()) @pytest.mark.parametrize( "shape", From aba43c3b22f24855d5931b8cde907db96c4e2757 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Sun, 27 Aug 2023 12:16:52 +0200 Subject: [PATCH 5/6] Fixed typos --- dpnp/dpnp_iface.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/dpnp/dpnp_iface.py b/dpnp/dpnp_iface.py index 5060ace272c4..69195b9c3f50 100644 --- a/dpnp/dpnp_iface.py +++ b/dpnp/dpnp_iface.py @@ -463,12 +463,12 @@ def is_supported_array_or_scalar(a): Parameters ---------- a : array - A input scalar or an array to check the type of. + An input scalar or an array to check the type of. Returns ------- out : bool - ``True`` if input `a` is a scalar or an array of supportedtype, + ``True`` if input `a` is a scalar or an array of supported type, ``False`` otherwise. """ From e82fa5698bcce3ead480038224102062206dac65 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Mon, 28 Aug 2023 13:21:05 +0200 Subject: [PATCH 6/6] Addressed review commets --- .github/workflows/conda-package.yml | 1 + dpnp/dpnp_iface.py | 4 ++-- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/.github/workflows/conda-package.yml b/.github/workflows/conda-package.yml index ebb1697ae378..2625952b5809 100644 --- a/.github/workflows/conda-package.yml +++ b/.github/workflows/conda-package.yml @@ -23,6 +23,7 @@ env: test_random_state.py test_sort.py test_special.py + test_sycl_queue.py test_umath.py test_usm_type.py third_party/cupy/linalg_tests/test_product.py diff --git a/dpnp/dpnp_iface.py b/dpnp/dpnp_iface.py index 69195b9c3f50..c131bf733ce1 100644 --- a/dpnp/dpnp_iface.py +++ b/dpnp/dpnp_iface.py @@ -462,7 +462,7 @@ def is_supported_array_or_scalar(a): Parameters ---------- - a : array + a : {scalar, dpnp_array, usm_ndarray} An input scalar or an array to check the type of. Returns @@ -483,7 +483,7 @@ def is_supported_array_type(a): Parameters ---------- - a : array + a : {dpnp_array, usm_ndarray} An input array to check the type. Returns