ENH: update random.logistic backend; kernels (#666)

Author: samir-nasibli

dpnp/backend/kernels/dpnp_krnl_random.cpp

@@ -394,30 +394,38 @@ void dpnp_rng_laplace_c(void* result, const double loc, const double scale, cons
     event_out.wait();
 }
 
+template <typename _KernelNameSpecialization>
+class dpnp_rng_logistic_c_kernel;
+
 /*   Logistic(loc, scale) ~ loc + scale * log(u/(1.0 - u)) */
 template <typename _DataType>
 void dpnp_rng_logistic_c(void* result, const double loc, const double scale, const size_t size)
 {
-    if (!size)
+    if (!size || !result)
     {
         return;
     }
-    cl::sycl::vector_class<cl::sycl::event> no_deps;
 
     const _DataType d_zero = _DataType(0.0);
     const _DataType d_one = _DataType(1.0);
 
     _DataType* result1 = reinterpret_cast<_DataType*>(result);
 
     mkl_rng::uniform<_DataType> distribution(d_zero, d_one);
-    auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1);
-    event_out.wait();
-
-    for (size_t i = 0; i < size; i++)
-        result1[i] = log(result1[i] / (1.0 - result1[i]));
+    auto event_distribution = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1);
 
-    for (size_t i = 0; i < size; i++)
+    cl::sycl::range<1> gws(size);
+    auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
+        size_t i = global_id[0];
+        result1[i] = cl::sycl::log(result1[i] / (1.0 - result1[i]));
         result1[i] = loc + scale * result1[i];
+    };
+    auto kernel_func = [&](cl::sycl::handler& cgh) {
+        cgh.depends_on({event_distribution});
+        cgh.parallel_for<class dpnp_rng_logistic_c_kernel<_DataType>>(gws, kernel_parallel_for_func);
+    };
+    auto event = DPNP_QUEUE.submit(kernel_func);
+    event.wait();
 }
 
 template <typename _DataType>