From 1ec06993e1fa5a0f28edc05878bf2c627b6c889c Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Thu, 16 Feb 2023 08:21:48 -0600 Subject: [PATCH 1/2] Remove mixed host\dev implementation from dpnp.any() --- .gitignore | 4 +++ dpnp/backend/kernels/dpnp_krnl_logic.cpp | 43 +++++++++++++++++------- dpnp/dpnp_array.py | 13 ++++--- dpnp/dpnp_iface_logic.py | 37 +++++++++++--------- tests/test_arraycreation.py | 2 +- tests/test_logic.py | 2 +- tests/test_usm_type.py | 1 - 7 files changed, 67 insertions(+), 35 deletions(-) diff --git a/.gitignore b/.gitignore index fda4c1635310..7ed68aab8567 100644 --- a/.gitignore +++ b/.gitignore @@ -1,6 +1,7 @@ # CMake build and local install directory build build_cython +dpnp.egg-info # Byte-compiled / optimized / DLL files __pycache__/ @@ -14,6 +15,9 @@ coverage.xml # Backup files kept after git merge/rebase *.orig +# Build examples +example3 + *dpnp_backend* dpnp/**/*.cpython*.so dpnp/**/*.pyd diff --git a/dpnp/backend/kernels/dpnp_krnl_logic.cpp b/dpnp/backend/kernels/dpnp_krnl_logic.cpp index 157347aa90c0..e0e5baf733a9 100644 --- a/dpnp/backend/kernels/dpnp_krnl_logic.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_logic.cpp @@ -218,6 +218,8 @@ DPCTLSyclEventRef dpnp_any_c(DPCTLSyclQueueRef q_ref, 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; @@ -229,38 +231,50 @@ DPCTLSyclEventRef dpnp_any_c(DPCTLSyclQueueRef q_ref, } sycl::queue q = *(reinterpret_cast(q_ref)); - sycl::event event; - DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, size); - DPNPC_ptr_adapter<_ResultType> result1_ptr(q_ref, result1, 1, true, true); - const _DataType* array_in = input1_ptr.get_ptr(); - _ResultType* result = result1_ptr.get_ptr(); + const _DataType* array_in = static_cast(array1_in); + bool* result = static_cast(result1); - result[0] = false; + auto fill_event = q.fill(result, false, 1); if (!size) { - return event_ref; + event_ref = reinterpret_cast(&fill_event); + return DPCTLEvent_Copy(event_ref); } - sycl::range<1> gws(size); - auto kernel_parallel_for_func = [=](sycl::id<1> global_id) { - size_t i = global_id[0]; + 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 sg = nd_it.get_sub_group(); + const auto max_sg_size = sg.get_max_local_range()[0]; + const size_t start = + vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + sg.get_group_id()[0] * max_sg_size); + const size_t end = sycl::min(start + vec_sz * max_sg_size, size); - if (array_in[i]) + // each work-item reduces over "vec_sz" elements in the input array + bool local_reduction = sycl::joint_any_of( + sg, &array_in[start], &array_in[end], [&](_DataType elem) { return elem != static_cast<_DataType>(0); }); + + if (sg.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); }; - event = q.submit(kernel_func); + auto event = q.submit(kernel_func); event_ref = reinterpret_cast(&event); - return DPCTLEvent_Copy(event_ref); } @@ -275,6 +289,7 @@ void dpnp_any_c(const void* array1_in, void* result1, const size_t size) size, dep_event_vec_ref); DPCTLEvent_WaitAndThrow(event_ref); + DPCTLEvent_Delete(event_ref); } template @@ -829,6 +844,8 @@ void func_map_init_logic(func_map_t& fmap) fmap[DPNPFuncName::DPNP_FN_ANY_EXT][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_any_ext_c}; 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}; + fmap[DPNPFuncName::DPNP_FN_ANY_EXT][eft_C64][eft_C64] = {eft_C64, (void*)dpnp_any_ext_c, bool>}; + fmap[DPNPFuncName::DPNP_FN_ANY_EXT][eft_C128][eft_C128] = {eft_C128, (void*)dpnp_any_ext_c, bool>}; func_map_logic_1arg_1type_helper(fmap); func_map_logic_2arg_2type_helper(fmap); diff --git a/dpnp/dpnp_array.py b/dpnp/dpnp_array.py index c50ed9792720..29917081b5fc 100644 --- a/dpnp/dpnp_array.py +++ b/dpnp/dpnp_array.py @@ -345,19 +345,24 @@ def all(self, axis=None, out=None, keepdims=False): return dpnp.all(self, axis, out, keepdims) - def any(self, axis=None, out=None, keepdims=False): + def any(self, + axis=None, + out=None, + keepdims=False, + *, + where=True): """ Returns True if any of the elements of `a` evaluate to True. - Refer to `numpy.any` for full documentation. + Refer to :obj:`dpnp.any` for full documentation. See Also -------- - :obj:`numpy.any` : equivalent function + :obj:`dpnp.any` : equivalent function """ - return dpnp.any(self, axis, out, keepdims) + return dpnp.any(self, axis=axis, out=out, keepdims=keepdims, where=where) def argmax(self, axis=None, out=None): """ diff --git a/dpnp/dpnp_iface_logic.py b/dpnp/dpnp_iface_logic.py index e94b0f6c1efb..2f37f054a84c 100644 --- a/dpnp/dpnp_iface_logic.py +++ b/dpnp/dpnp_iface_logic.py @@ -163,7 +163,13 @@ def allclose(x1, x2, rtol=1.e-5, atol=1.e-8, **kwargs): return call_origin(numpy.allclose, x1, x2, rtol=rtol, atol=atol, **kwargs) -def any(x1, axis=None, out=None, keepdims=False): +def any(x1, + /, + axis=None, + out=None, + keepdims=False, + *, + where=True): """ Test whether any array element along a given axis evaluates to True. @@ -174,9 +180,10 @@ def any(x1, axis=None, out=None, keepdims=False): Input array is supported as :obj:`dpnp.ndarray`. Otherwise the function will be executed sequentially on CPU. Input array data types are limited by supported DPNP :ref:`Data types`. - Parameter ``axis`` is supported only with default value ``None``. - Parameter ``out`` is supported only with default value ``None``. - Parameter ``keepdims`` is supported only with default value ``False``. + Parameter `axis` is supported only with default value `None`. + Parameter `out` is supported only with default value `None`. + Parameter `keepdims` is supported only with default value `False`. + Parameter `where` is supported only with default value `True`. See Also -------- @@ -189,15 +196,15 @@ def any(x1, axis=None, out=None, keepdims=False): Examples -------- - >>> import dpnp as np - >>> x = np.array([[True, False], [True, True]]) - >>> np.any(x) + >>> import dpnp as dp + >>> x = dp.array([[True, False], [True, True]]) + >>> dp.any(x) True - >>> x2 = np.array([0, 0, 0]) - >>> np.any(x2) + >>> x2 = dp.array([0, 0, 0]) + >>> dp.any(x2) False - >>> x3 = np.array([1.0, np.nan]) - >>> np.any(x3) + >>> x3 = dp.array([1.0, dp.nan]) + >>> dp.any(x3) True """ @@ -210,13 +217,13 @@ def any(x1, axis=None, out=None, keepdims=False): pass elif keepdims is not False: pass + elif where is not True: + pass else: result_obj = dpnp_any(x1_desc).get_pyobj() - result = dpnp.convert_single_elem_array_to_scalar(result_obj) - - return result + return dpnp.convert_single_elem_array_to_scalar(result_obj) - return call_origin(numpy.any, x1, axis, out, keepdims) + return call_origin(numpy.any, x1, axis=axis, out=out, keepdims=keepdims, where=where) def equal(x1, diff --git a/tests/test_arraycreation.py b/tests/test_arraycreation.py index 63435bca11f0..71c8e8bcb812 100644 --- a/tests/test_arraycreation.py +++ b/tests/test_arraycreation.py @@ -507,7 +507,7 @@ def test_dpctl_tensor_input(func, args): new_args = [eval(val, {'x0' : x0}) for val in args] X = getattr(dpt, func)(*new_args) Y = getattr(dpnp, func)(*new_args) - if func is 'empty_like': + if func == 'empty_like': assert X.shape == Y.shape else: assert_array_equal(X, Y) diff --git a/tests/test_logic.py b/tests/test_logic.py index 425106fd2efe..17e07dd613d6 100644 --- a/tests/test_logic.py +++ b/tests/test_logic.py @@ -63,7 +63,7 @@ def test_allclose(type): assert_allclose(dpnp_res, np_res) -@pytest.mark.parametrize("type", get_all_dtypes(no_complex=True)) +@pytest.mark.parametrize("type", get_all_dtypes()) @pytest.mark.parametrize("shape", [(0,), (4,), (2, 3), (2, 2, 2)], ids=['(0,)', '(4,)', '(2,3)', '(2,2,2)']) diff --git a/tests/test_usm_type.py b/tests/test_usm_type.py index 1a33a1d655dd..008a296a1b7b 100644 --- a/tests/test_usm_type.py +++ b/tests/test_usm_type.py @@ -90,7 +90,6 @@ def test_array_creation(func, args, usm_type_x, usm_type_y): assert y.usm_type == usm_type_y -@pytest.mark.skip() @pytest.mark.parametrize("func", ["tril", "triu"], ids=["tril", "triu"]) @pytest.mark.parametrize("usm_type", list_of_usm_types, ids=list_of_usm_types) def test_tril_triu(func, usm_type): From 95581ee2f1e3cc046720744b26b9904761ef2310 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Fri, 24 Feb 2023 04:55:29 -0600 Subject: [PATCH 2/2] Reduce over group --- dpnp/backend/kernels/dpnp_krnl_logic.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/dpnp/backend/kernels/dpnp_krnl_logic.cpp b/dpnp/backend/kernels/dpnp_krnl_logic.cpp index e0e5baf733a9..d92a3c5deb85 100644 --- a/dpnp/backend/kernels/dpnp_krnl_logic.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_logic.cpp @@ -251,17 +251,17 @@ 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 sg = nd_it.get_sub_group(); - const auto max_sg_size = sg.get_max_local_range()[0]; + auto gr = nd_it.get_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) + sg.get_group_id()[0] * max_sg_size); - const size_t end = sycl::min(start + vec_sz * max_sg_size, size); + 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( - sg, &array_in[start], &array_in[end], [&](_DataType elem) { return elem != static_cast<_DataType>(0); }); + gr, &array_in[start], &array_in[end], [&](_DataType elem) { return elem != static_cast<_DataType>(0); }); - if (sg.leader() && (local_reduction == true)) + if (gr.leader() && (local_reduction == true)) { result[0] = true; }