Enable Compute Follows Data in Cython

dpnp/backend/include/dpnp_gen_1arg_1type_tbl.hpp

@@ -90,20 +90,20 @@
 
 #endif
 
-MACRO_1ARG_1TYPE_OP(dpnp_conjugate_c, std::conj(input_elem), DPNP_QUEUE.submit(kernel_func))
-MACRO_1ARG_1TYPE_OP(dpnp_copy_c, input_elem, DPNP_QUEUE.submit(kernel_func))
+MACRO_1ARG_1TYPE_OP(dpnp_conjugate_c, std::conj(input_elem), q.submit(kernel_func))
+MACRO_1ARG_1TYPE_OP(dpnp_copy_c, input_elem, q.submit(kernel_func))
 MACRO_1ARG_1TYPE_OP(dpnp_erf_c,
                     sycl::erf((double)input_elem),
-                    oneapi::mkl::vm::erf(DPNP_QUEUE, input1_size, input1_data, result)) // no sycl::erf for int and long
-MACRO_1ARG_1TYPE_OP(dpnp_negative_c, -input_elem, DPNP_QUEUE.submit(kernel_func))
+                    oneapi::mkl::vm::erf(q, input1_size, input1_data, result)) // no sycl::erf for int and long
+MACRO_1ARG_1TYPE_OP(dpnp_negative_c, -input_elem, q.submit(kernel_func))
 MACRO_1ARG_1TYPE_OP(dpnp_recip_c,
                     _DataType(1) / input_elem,
-                    DPNP_QUEUE.submit(kernel_func)) // error: no member named 'recip' in namespace 'sycl'
+                    q.submit(kernel_func)) // error: no member named 'recip' in namespace 'sycl'
 MACRO_1ARG_1TYPE_OP(dpnp_sign_c,
                     sycl::sign((double)input_elem),
-                    DPNP_QUEUE.submit(kernel_func)) // no sycl::sign for int and long
+                    q.submit(kernel_func)) // no sycl::sign for int and long
 MACRO_1ARG_1TYPE_OP(dpnp_square_c,
                     input_elem* input_elem,
-                    oneapi::mkl::vm::sqr(DPNP_QUEUE, input1_size, input1_data, result))
+                    oneapi::mkl::vm::sqr(q, input1_size, input1_data, result))
 
 #undef MACRO_1ARG_1TYPE_OP

dpnp/backend/include/dpnp_gen_1arg_2type_tbl.hpp

@@ -92,81 +92,81 @@
 
 MACRO_1ARG_2TYPES_OP(dpnp_acos_c,
                      sycl::acos(input_elem),
-                     oneapi::mkl::vm::acos(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::acos(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_acosh_c,
                      sycl::acosh(input_elem),
-                     oneapi::mkl::vm::acosh(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::acosh(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_asin_c,
                      sycl::asin(input_elem),
-                     oneapi::mkl::vm::asin(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::asin(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_asinh_c,
                      sycl::asinh(input_elem),
-                     oneapi::mkl::vm::asinh(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::asinh(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_atan_c,
                      sycl::atan(input_elem),
-                     oneapi::mkl::vm::atan(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::atan(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_atanh_c,
                      sycl::atanh(input_elem),
-                     oneapi::mkl::vm::atanh(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::atanh(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_cbrt_c,
                      sycl::cbrt(input_elem),
-                     oneapi::mkl::vm::cbrt(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::cbrt(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_ceil_c,
                      sycl::ceil(input_elem),
-                     oneapi::mkl::vm::ceil(DPNP_QUEUE, input1_size, input1_data, result))
-MACRO_1ARG_2TYPES_OP(dpnp_copyto_c, input_elem, DPNP_QUEUE.submit(kernel_func))
+                     oneapi::mkl::vm::ceil(q, input1_size, input1_data, result))
+MACRO_1ARG_2TYPES_OP(dpnp_copyto_c, input_elem, q.submit(kernel_func))
 MACRO_1ARG_2TYPES_OP(dpnp_cos_c,
                      sycl::cos(input_elem),
-                     oneapi::mkl::vm::cos(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::cos(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_cosh_c,
                      sycl::cosh(input_elem),
-                     oneapi::mkl::vm::cosh(DPNP_QUEUE, input1_size, input1_data, result))
-MACRO_1ARG_2TYPES_OP(dpnp_degrees_c, sycl::degrees(input_elem), DPNP_QUEUE.submit(kernel_func))
+                     oneapi::mkl::vm::cosh(q, input1_size, input1_data, result))
+MACRO_1ARG_2TYPES_OP(dpnp_degrees_c, sycl::degrees(input_elem), q.submit(kernel_func))
 MACRO_1ARG_2TYPES_OP(dpnp_exp2_c,
                      sycl::exp2(input_elem),
-                     oneapi::mkl::vm::exp2(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::exp2(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_exp_c,
                      sycl::exp(input_elem),
-                     oneapi::mkl::vm::exp(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::exp(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_expm1_c,
                      sycl::expm1(input_elem),
-                     oneapi::mkl::vm::expm1(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::expm1(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_fabs_c,
                      sycl::fabs(input_elem),
-                     oneapi::mkl::vm::abs(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::abs(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_floor_c,
                      sycl::floor(input_elem),
-                     oneapi::mkl::vm::floor(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::floor(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_log10_c,
                      sycl::log10(input_elem),
-                     oneapi::mkl::vm::log10(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::log10(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_log1p_c,
                      sycl::log1p(input_elem),
-                     oneapi::mkl::vm::log1p(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::log1p(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_log2_c,
                      sycl::log2(input_elem),
-                     oneapi::mkl::vm::log2(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::log2(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_log_c,
                      sycl::log(input_elem),
-                     oneapi::mkl::vm::ln(DPNP_QUEUE, input1_size, input1_data, result))
-MACRO_1ARG_2TYPES_OP(dpnp_radians_c, sycl::radians(input_elem), DPNP_QUEUE.submit(kernel_func))
+                     oneapi::mkl::vm::ln(q, input1_size, input1_data, result))
+MACRO_1ARG_2TYPES_OP(dpnp_radians_c, sycl::radians(input_elem), q.submit(kernel_func))
 MACRO_1ARG_2TYPES_OP(dpnp_sin_c,
                      sycl::sin(input_elem),
-                     oneapi::mkl::vm::sin(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::sin(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_sinh_c,
                      sycl::sinh(input_elem),
-                     oneapi::mkl::vm::sinh(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::sinh(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_sqrt_c,
                      sycl::sqrt(input_elem),
-                     oneapi::mkl::vm::sqrt(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::sqrt(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_tan_c,
                      sycl::tan(input_elem),
-                     oneapi::mkl::vm::tan(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::tan(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_tanh_c,
                      sycl::tanh(input_elem),
-                     oneapi::mkl::vm::tanh(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::tanh(q, input1_size, input1_data, result))
 MACRO_1ARG_2TYPES_OP(dpnp_trunc_c,
                      sycl::trunc(input_elem),
-                     oneapi::mkl::vm::trunc(DPNP_QUEUE, input1_size, input1_data, result))
+                     oneapi::mkl::vm::trunc(q, input1_size, input1_data, result))
 
 #undef MACRO_1ARG_2TYPES_OP

dpnp/backend/include/dpnp_iface.hpp

@@ -533,7 +533,7 @@ INP_DLLEXPORT void dpnp_cumprod_c(void* array1_in, void* result1, size_t size);
  * @param [in]  q_ref               Reference to SYCL queue.
  * @param [in]  array1_in           Input array.
  * @param [out] result1             Output array.
- * @param [in]  size                 Number of elements in input arrays.
+ * @param [in]  size                Number of elements in input arrays.
  * @param [in]  dep_event_vec_ref   Reference to vector of SYCL events.
  *
  */
@@ -551,7 +551,7 @@ INP_DLLEXPORT void dpnp_cumsum_c(void* array1_in, void* result1, size_t size);
  * @ingroup BACKEND_API
  * @brief The differences between consecutive elements of an array.
  *
- * @param [in]  q_ref             Reference to SYCL queue.
+ * @param [in]  q_ref               Reference to SYCL queue.
  * @param [out] result_out          Output array.
  * @param [in]  result_size         Size of output array.
  * @param [in]  result_ndim         Number of output array dimensions.

dpnp/backend/include/dpnp_iface_fptr.hpp

@@ -381,7 +381,6 @@ enum class DPNPFuncName : size_t
     DPNP_FN_ZEROS_LIKE,                   /**< Used in numpy.zeros_like() impl */
     DPNP_FN_ZEROS_LIKE_EXT,               /**< Used in numpy.zeros_like() impl, requires extra parameters */
     DPNP_FN_LAST,                         /**< The latest element of the enumeration */
-    DPNP_FN_LAST_EXT                      /**< The latest element of the enumeration, requires extra parameters */
 };
 
 /**

dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp

@@ -73,9 +73,9 @@ DPCTLSyclEventRef dpnp_arange_c(DPCTLSyclQueueRef q_ref,
 
     event = q.submit(kernel_func);
 
-    event.wait();
+    event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);
 
-    return event_ref;
+    return DPCTLEvent_Copy(event_ref);
 }
 
 template <typename _DataType>
@@ -617,9 +617,6 @@ DPCTLSyclEventRef dpnp_vander_c(DPCTLSyclQueueRef q_ref,
                                 const int increasing,
                                 const DPCTLEventVectorRef dep_event_vec_ref)
 {
-    // avoid warning unused variable
-    (void)dep_event_vec_ref;
-
     DPCTLSyclEventRef event_ref = nullptr;
 
     if ((array1_in == nullptr) || (result1 == nullptr))
@@ -637,8 +634,7 @@ DPCTLSyclEventRef dpnp_vander_c(DPCTLSyclQueueRef q_ref,
 
     if (N == 1)
     {
-        dpnp_ones_c<_DataType_output>(result, size_in);
-        return event_ref;
+        return dpnp_ones_c<_DataType_output>(q_ref, result, size_in, dep_event_vec_ref);
     }
 
     if (increasing)
@@ -762,9 +758,10 @@ DPCTLSyclEventRef dpnp_trace_c(DPCTLSyclQueueRef q_ref,
     };
 
     auto event = q.submit(kernel_func);
-    event.wait();
 
-    return event_ref;
+    event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);
+
+    return DPCTLEvent_Copy(event_ref);
 }
 
 template <typename _DataType, typename _ResultType>
@@ -849,9 +846,9 @@ DPCTLSyclEventRef dpnp_tri_c(DPCTLSyclQueueRef q_ref,
 
     event = q.submit(kernel_func);
 
-    event.wait();
+    event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);
 
-    return event_ref;
+    return DPCTLEvent_Copy(event_ref);
 }
 
 template <typename _DataType>

dpnp/backend/kernels/dpnp_krnl_bitwise.cpp

@@ -68,9 +68,9 @@ DPCTLSyclEventRef dpnp_invert_c(DPCTLSyclQueueRef q_ref,
 
     event = q.submit(kernel_func);
 
-    event.wait();
+    event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);
 
-    return event_ref;
+    return DPCTLEvent_Copy(event_ref);
 }
 
 template <typename _DataType>
@@ -101,6 +101,9 @@ static void func_map_init_bitwise_1arg_1type(func_map_t& fmap)
     fmap[DPNPFuncName::DPNP_FN_INVERT][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_invert_default_c<int32_t>};
     fmap[DPNPFuncName::DPNP_FN_INVERT][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_invert_default_c<int64_t>};
 
+    fmap[DPNPFuncName::DPNP_FN_INVERT_EXT][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_invert_ext_c<int32_t>};
+    fmap[DPNPFuncName::DPNP_FN_INVERT_EXT][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_invert_ext_c<int64_t>};
+
     return;
 }
 
@@ -209,7 +212,6 @@ static void func_map_init_bitwise_1arg_1type(func_map_t& fmap)
                 cgh.parallel_for<class __name__##_strides_kernel<_DataType>>(gws, kernel_parallel_for_func);           \
             };                                                                                                         \
             event = q.submit(kernel_func);                                                                             \
-            event.wait();                                                                                              \
         }                                                                                                              \
         else                                                                                                           \
         {                                                                                                              \
@@ -223,9 +225,10 @@ static void func_map_init_bitwise_1arg_1type(func_map_t& fmap)
                 cgh.parallel_for<class __name__##_kernel<_DataType>>(gws, kernel_parallel_for_func);                   \
             };                                                                                                         \
             event = q.submit(kernel_func);                                                                             \
-            event.wait();                                                                                              \
         }                                                                                                              \
-        return event_ref;                                                                             \
+        event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);                                                       \
+                                                                                                                       \
+        return DPCTLEvent_Copy(event_ref);                                                                             \
     }                                                                                                                  \
                                                                                                                        \
     template <typename _DataType>                                                                                      \

dpnp/backend/kernels/dpnp_krnl_common.cpp

@@ -81,9 +81,9 @@ DPCTLSyclEventRef dpnp_astype_c(DPCTLSyclQueueRef q_ref,
 
     event = q.submit(kernel_func);
 
-    event.wait();
+    event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);
 
-    return event_ref;
+    return DPCTLEvent_Copy(event_ref);
 }
 
 template <typename _DataType, typename _ResultType>
@@ -732,9 +732,9 @@ DPCTLSyclEventRef dpnp_initval_c(DPCTLSyclQueueRef q_ref,
 
     sycl::event event = q.submit(kernel_func);
 
-    event.wait();
+    event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);
 
-    return event_ref;
+    return DPCTLEvent_Copy(event_ref);
 }
 
 template <typename _DataType>
@@ -756,10 +756,10 @@ void (*dpnp_initval_default_c)(void*, void*, size_t) = dpnp_initval_c<_DataType>
 
 template <typename _DataType>
 DPCTLSyclEventRef (*dpnp_initval_ext_c)(DPCTLSyclQueueRef,
-                                    void*,
-                                    void*,
-                                    size_t,
-                                    const DPCTLEventVectorRef) = dpnp_initval_c<_DataType>;
+                                        void*,
+                                        void*,
+                                        size_t,
+                                        const DPCTLEventVectorRef) = dpnp_initval_c<_DataType>;
 
 template <typename _KernelNameSpecialization>
 class dpnp_matmul_c_kernel;