ELEMWISE 1arg_1types add 1 MKL kernel (#137)

* ELEMWISE 1arg_1types add MKL kernels

Author: densmirn

dpnp/backend/backend_iface.hpp

@@ -396,7 +396,7 @@ INP_DLLEXPORT void custom_var_c(
  * @param [in]  size  Number of elements in the input array.
  *
  */
-#define MACRO_CUSTOM_1ARG_1TYPE_OP(__name__, __operation__)                                                            \
+#define MACRO_CUSTOM_1ARG_1TYPE_OP(__name__, __operation1__, __operation2__)                                           \
     template <typename _DataType>                                                                                      \
     INP_DLLEXPORT void __name__(void* array1, void* result1, size_t size);
 

dpnp/backend/custom_1arg_1type_tbl.hpp

@@ -31,16 +31,19 @@
  * Parameters:
  * - public name of the function and kernel name
  * - operation used to calculate the result
+ * - mkl operation used to calculate the result
  *
  */
 
 #ifndef MACRO_CUSTOM_1ARG_1TYPE_OP
 #error "MACRO_CUSTOM_1ARG_1TYPE_OP is not defined"
 #endif
 
-MACRO_CUSTOM_1ARG_1TYPE_OP(dpnp_recip_c,
-                           _DataType(1) / input_elem) // error: no member named 'recip' in namespace 'cl::sycl'
-MACRO_CUSTOM_1ARG_1TYPE_OP(dpnp_sign_c, cl::sycl::sign((double)input_elem)) // no sycl::sign for int and long
-MACRO_CUSTOM_1ARG_1TYPE_OP(dpnp_square_c, input_elem* input_elem)
+MACRO_CUSTOM_1ARG_1TYPE_OP(dpnp_recip_c, _DataType(1) / input_elem,
+                           DPNP_QUEUE.submit(kernel_func)) // error: no member named 'recip' in namespace 'cl::sycl'
+MACRO_CUSTOM_1ARG_1TYPE_OP(dpnp_sign_c, cl::sycl::sign((double)input_elem),
+                           DPNP_QUEUE.submit(kernel_func)) // no sycl::sign for int and long
+MACRO_CUSTOM_1ARG_1TYPE_OP(dpnp_square_c, input_elem * input_elem,
+                           oneapi::mkl::vm::sqr(DPNP_QUEUE, size, array1, result))
 
 #undef MACRO_CUSTOM_1ARG_1TYPE_OP

dpnp/backend/custom_kernels_elemwise.cpp

@@ -202,14 +202,15 @@ static void func_map_init_elemwise_1arg_2type(func_map_t& fmap)
     return;
 }
 
-#define MACRO_CUSTOM_1ARG_1TYPE_OP(__name__, __operation__)                                                            \
+#define MACRO_CUSTOM_1ARG_1TYPE_OP(__name__, __operation1__, __operation2__)                                           \
     template <typename _KernelNameSpecialization>                                                                      \
     class __name__##_kernel;                                                                                           \
                                                                                                                        \
     template <typename _DataType>                                                                                      \
     void __name__(void* array1_in, void* result1, size_t size)                                                         \
     {                                                                                                                  \
         cl::sycl::event event;                                                                                         \
+                                                                                                                       \
         _DataType* array1 = reinterpret_cast<_DataType*>(array1_in);                                                   \
         _DataType* result = reinterpret_cast<_DataType*>(result1);                                                     \
                                                                                                                        \
@@ -218,15 +219,22 @@ static void func_map_init_elemwise_1arg_2type(func_map_t& fmap)
             size_t i = global_id[0]; /*for (size_t i = 0; i < size; ++i)*/                                             \
             {                                                                                                          \
                 _DataType input_elem = array1[i];                                                                      \
-                result[i] = __operation__;                                                                             \
+                result[i] = __operation1__;                                                                            \
             }                                                                                                          \
         };                                                                                                             \
                                                                                                                        \
         auto kernel_func = [&](cl::sycl::handler& cgh) {                                                               \
             cgh.parallel_for<class __name__##_kernel<_DataType>>(gws, kernel_parallel_for_func);                       \
         };                                                                                                             \
                                                                                                                        \
-        event = DPNP_QUEUE.submit(kernel_func);                                                                        \
+        if constexpr (std::is_same<_DataType, double>::value)                                                          \
+        {                                                                                                              \
+            event = __operation2__;                                                                                    \
+        }                                                                                                              \
+        else                                                                                                           \
+        {                                                                                                              \
+            event = DPNP_QUEUE.submit(kernel_func);                                                                    \
+        }                                                                                                              \
                                                                                                                        \
         event.wait();                                                                                                  \
     }