Remove mixed host\dev implementation from dpnp.any()

Author: antonwolfy

.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

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<sycl::queue*>(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<const _DataType*>(array1_in);
+    bool* result = static_cast<bool*>(result1);
 
-    result[0] = false;
+    auto fill_event = q.fill(result, false, 1);
 
     if (!size)
     {
-        return event_ref;
+        event_ref = reinterpret_cast<DPCTLSyclEventRef>(&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<class dpnp_any_c_kernel<_DataType, _ResultType>>(gws, kernel_parallel_for_func);
     };
 
-    event = q.submit(kernel_func);
+    auto event = q.submit(kernel_func);
 
     event_ref = reinterpret_cast<DPCTLSyclEventRef>(&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 <typename _DataType, typename _ResultType>
@@ -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<int64_t, bool>};
     fmap[DPNPFuncName::DPNP_FN_ANY_EXT][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_any_ext_c<float, bool>};
     fmap[DPNPFuncName::DPNP_FN_ANY_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_any_ext_c<double, bool>};
+    fmap[DPNPFuncName::DPNP_FN_ANY_EXT][eft_C64][eft_C64] = {eft_C64, (void*)dpnp_any_ext_c<std::complex<float>, bool>};
+    fmap[DPNPFuncName::DPNP_FN_ANY_EXT][eft_C128][eft_C128] = {eft_C128, (void*)dpnp_any_ext_c<std::complex<double>, bool>};
 
     func_map_logic_1arg_1type_helper<eft_BLN, eft_INT, eft_LNG, eft_FLT, eft_DBL>(fmap);
     func_map_logic_2arg_2type_helper<eft_BLN, eft_INT, eft_LNG, eft_FLT, eft_DBL>(fmap);

dpnp/dpnp_array.py

@@ -327,19 +327,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):
         """

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,

tests/test_arraycreation.py

@@ -505,7 +505,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)

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)'])

tests/test_usm_type.py

@@ -64,7 +64,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):