NEGATIVE add kernel required for black_scholes (#674)

Author: densmirn

dpnp/backend/include/dpnp_gen_1arg_1type_tbl.hpp

@@ -57,7 +57,10 @@
 
 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_erf_c, input_elem, oneapi::mkl::vm::erf(DPNP_QUEUE, size, array1, result))
+MACRO_1ARG_1TYPE_OP(dpnp_erf_c,
+                    cl::sycl::erf((double)input_elem),
+                    oneapi::mkl::vm::erf(DPNP_QUEUE, size, array1, result)) // no sycl::erf for int and long
+MACRO_1ARG_1TYPE_OP(dpnp_negative_c, -input_elem, DPNP_QUEUE.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 'cl::sycl'

dpnp/backend/include/dpnp_iface_fptr.hpp

@@ -134,6 +134,7 @@ enum class DPNPFuncName : size_t
     DPNP_FN_MINIMUM,                  /**< Used in numpy.minimum() implementation  */
     DPNP_FN_MODF,                     /**< Used in numpy.modf() implementation  */
     DPNP_FN_MULTIPLY,                 /**< Used in numpy.multiply() implementation  */
+    DPNP_FN_NEGATIVE,                 /**< Used in numpy.negative() implementation  */
     DPNP_FN_NONZERO,                  /**< Used in numpy.nonzero() implementation  */
     DPNP_FN_ONES,                     /**< Used in numpy.ones() implementation */
     DPNP_FN_ONES_LIKE,                /**< Used in numpy.ones_like() implementation */

dpnp/backend/kernels/dpnp_krnl_elemwise.cpp

@@ -317,6 +317,11 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap)
     fmap[DPNPFuncName::DPNP_FN_FLATTEN][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_copy_c<double>};
     fmap[DPNPFuncName::DPNP_FN_FLATTEN][eft_C128][eft_C128] = {eft_C128, (void*)dpnp_copy_c<std::complex<double>>};
 
+    fmap[DPNPFuncName::DPNP_FN_NEGATIVE][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_negative_c<int>};
+    fmap[DPNPFuncName::DPNP_FN_NEGATIVE][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_negative_c<long>};
+    fmap[DPNPFuncName::DPNP_FN_NEGATIVE][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_negative_c<float>};
+    fmap[DPNPFuncName::DPNP_FN_NEGATIVE][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_negative_c<double>};
+
     fmap[DPNPFuncName::DPNP_FN_RECIP][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_recip_c<int>};
     fmap[DPNPFuncName::DPNP_FN_RECIP][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_recip_c<long>};
     fmap[DPNPFuncName::DPNP_FN_RECIP][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_recip_c<float>};

dpnp/dpnp_algo/dpnp_algo.pxd

@@ -107,6 +107,7 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName":  # need this na
         DPNP_FN_MINIMUM
         DPNP_FN_MODF
         DPNP_FN_MULTIPLY
+        DPNP_FN_NEGATIVE
         DPNP_FN_NONZERO
         DPNP_FN_ONES
         DPNP_FN_ONES_LIKE

dpnp/dpnp_algo/dpnp_algo_mathematical.pyx

@@ -330,12 +330,7 @@ cpdef dpnp_nansum(dparray x1):
 
 
 cpdef dparray dpnp_negative(dparray array1):
-    cdef dparray result = dparray(array1.shape, dtype=array1.dtype)
-
-    for i in range(result.size):
-        result[i] = -(array1[i])
-
-    return result
+    return call_fptr_1in_1out(DPNP_FN_NEGATIVE, array1, array1.shape)
 
 
 cpdef dparray dpnp_power(object x1_obj, object x2_obj, object dtype=None, dparray out=None, object where=True):

tests/test_mathematical.py

@@ -197,6 +197,21 @@ def test_nancumsum(array):
     numpy.testing.assert_array_equal(expected, result)
 
 
+@pytest.mark.parametrize("data",
+                         [[[1., -1.], [0.1, -0.1]], [-2, -1, 0, 1, 2]],
+                         ids=['[[1., -1.], [0.1, -0.1]]', '[-2, -1, 0, 1, 2]'])
+@pytest.mark.parametrize("dtype",
+                         [numpy.float64, numpy.float32, numpy.int64, numpy.int32],
+                         ids=['numpy.float64', 'numpy.float32', 'numpy.int64', 'numpy.int32'])
+def test_negative(data, dtype):
+    a = numpy.array(data, dtype=dtype)
+    ia = inp.array(data, dtype=dtype)
+
+    result = inp.negative(ia)
+    expected = numpy.negative(a)
+    numpy.testing.assert_array_equal(result, expected)
+
+
 @pytest.mark.parametrize("val_type",
                          [numpy.float64, numpy.float32, numpy.int64, numpy.int32],
                          ids=['numpy.float64', 'numpy.float32', 'numpy.int64', 'numpy.int32'])