add kernel in functions all, any (#640)

* add kernel in functions all, any

Author: Rubtsowa

dpnp/backend/include/dpnp_iface.hpp

@@ -114,7 +114,7 @@ void dpnp_memory_memcpy_c(void* dst, const void* src, size_t size_in_bytes);
  * @param [in]  size        Number of input elements in `array`.
  */
 template <typename _DataType, typename _ResultType>
-INP_DLLEXPORT void dpnp_all_c(void* array, void* result, const size_t size);
+INP_DLLEXPORT void dpnp_all_c(const void* array, void* result, const size_t size);
 
 /**
  * @ingroup BACKEND_API

dpnp/backend/kernels/dpnp_krnl_logic.cpp

@@ -29,53 +29,90 @@
 #include "dpnp_iface.hpp"
 #include "queue_sycl.hpp"
 
+
+template <typename _DataType, typename _ResultType>
+class dpnp_all_c_kernel;
+
 template <typename _DataType, typename _ResultType>
-void dpnp_all_c(void* array1_in, void* result1, const size_t size)
+void dpnp_all_c(const void* array1_in, void* result1, const size_t size)
 {
-    _DataType* array_in = reinterpret_cast<_DataType*>(array1_in);
+    cl::sycl::event event;
+
+    const _DataType* array_in = reinterpret_cast<const _DataType*>(array1_in);
     _ResultType* result = reinterpret_cast<_ResultType*>(result1);
 
-    bool res = true;
+    if (!array1_in || !result1)
+    {
+        return;
+    }
+
+    result[0] = true;
 
-    for (size_t i = 0; i < size; ++i)
+    if (!size)
     {
+        return;
+    }
+
+    cl::sycl::range<1> gws(size);
+    auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
+        size_t i = global_id[0];
+
         if (!array_in[i])
         {
-            res = false;
-            break;
+            result[0] = false;
         }
-    }
+    };
 
-    result[0] = res;
+    auto kernel_func = [&](cl::sycl::handler& cgh) {
+        cgh.parallel_for<class dpnp_all_c_kernel<_DataType, _ResultType>>(gws, kernel_parallel_for_func);
+    };
 
-    return;
+    event = DPNP_QUEUE.submit(kernel_func);
+
+    event.wait();
 }
 
+template <typename _DataType, typename _ResultType>
+class dpnp_any_c_kernel;
+
 template <typename _DataType, typename _ResultType>
 void dpnp_any_c(const void* array1_in, void* result1, const size_t size)
 {
-    if ((array1_in == nullptr) || (result1 == nullptr))
-    {
-        return;
-    }
+    cl::sycl::event event;
 
     const _DataType* array_in = reinterpret_cast<const _DataType*>(array1_in);
     _ResultType* result = reinterpret_cast<_ResultType*>(result1);
 
-    bool res = false;
 
-    for (size_t i = 0; i < size; ++i)
+    if (!array1_in || !result1)
+    {
+        return;
+    }
+
+    result[0] = false;
+
+    if (!size)
     {
+        return;
+    }
+
+    cl::sycl::range<1> gws(size);
+    auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
+        size_t i = global_id[0];
+
         if (array_in[i])
         {
-            res = true;
-            break;
+            result[0] = true;
         }
-    }
+    };
 
-    result[0] = res;
+    auto kernel_func = [&](cl::sycl::handler& cgh) {
+        cgh.parallel_for<class dpnp_any_c_kernel<_DataType, _ResultType>>(gws, kernel_parallel_for_func);
+    };
 
-    return;
+    event = DPNP_QUEUE.submit(kernel_func);
+
+    event.wait();
 }
 
 void func_map_init_logic(func_map_t& fmap)

tests/test_logic.py

@@ -9,8 +9,8 @@
                          [numpy.float64, numpy.float32, numpy.int64, numpy.int32, numpy.bool, numpy.bool_],
                          ids=['float64', 'float32', 'int64', 'int32', 'bool', 'bool_'])
 @pytest.mark.parametrize("shape",
-                         [(4,), (2, 3), (2, 2, 2)],
-                         ids=['(4,)', '(2,3)', '(2,2,2)'])
+                         [(0,), (4,), (2, 3), (2, 2, 2)],
+                         ids=['(0,)', '(4,)', '(2,3)', '(2,2,2)'])
 def test_all(type, shape):
     size = 1
     for i in range(len(shape)):
@@ -42,8 +42,8 @@ def test_all(type, shape):
                          [numpy.float64, numpy.float32, numpy.int64, numpy.int32, numpy.bool, numpy.bool_],
                          ids=['float64', 'float32', 'int64', 'int32', 'bool', 'bool_'])
 @pytest.mark.parametrize("shape",
-                         [(4,), (2, 3), (2, 2, 2)],
-                         ids=['(4,)', '(2,3)', '(2,2,2)'])
+                         [(0,), (4,), (2, 3), (2, 2, 2)],
+                         ids=['(0,)', '(4,)', '(2,3)', '(2,2,2)'])
 def test_any(type, shape):
     size = 1
     for i in range(len(shape)):