ERF add kernel (#546)

Author: densmirn

dpnp/backend/include/dpnp_gen_1arg_1type_tbl.hpp

@@ -57,6 +57,7 @@
 
 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_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

@@ -90,6 +90,7 @@ enum class DPNPFuncName : size_t
     DPNP_FN_DOT,                      /**< Used in numpy.dot() implementation  */
     DPNP_FN_EIG,                      /**< Used in numpy.linalg.eig() implementation  */
     DPNP_FN_EIGVALS,                  /**< Used in numpy.linalg.eigvals() implementation  */
+    DPNP_FN_ERF,                      /**< Used in scipy.special.erf implementation  */
     DPNP_FN_EXP,                      /**< Used in numpy.exp() implementation  */
     DPNP_FN_EXP2,                     /**< Used in numpy.exp2() implementation  */
     DPNP_FN_EXPM1,                    /**< Used in numpy.expm1() implementation  */

dpnp/backend/kernels/dpnp_krnl_elemwise.cpp

@@ -235,7 +235,7 @@ static void func_map_init_elemwise_1arg_2type(func_map_t& fmap)
             cgh.parallel_for<class __name__##_kernel<_DataType>>(gws, kernel_parallel_for_func);                       \
         };                                                                                                             \
                                                                                                                        \
-        if constexpr (std::is_same<_DataType, double>::value)                                                          \
+        if constexpr (std::is_same<_DataType, double>::value || std::is_same<_DataType, float>::value)                 \
         {                                                                                                              \
             event = __operation2__;                                                                                    \
         }                                                                                                              \
@@ -297,6 +297,11 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap)
     fmap[DPNPFuncName::DPNP_FN_CONJIGUATE][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_copy_c<int>};
     fmap[DPNPFuncName::DPNP_FN_CONJIGUATE][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_copy_c<long>};
 
+    fmap[DPNPFuncName::DPNP_FN_ERF][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_erf_c<int>};
+    fmap[DPNPFuncName::DPNP_FN_ERF][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_erf_c<long>};
+    fmap[DPNPFuncName::DPNP_FN_ERF][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_erf_c<float>};
+    fmap[DPNPFuncName::DPNP_FN_ERF][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_erf_c<double>};
+
     fmap[DPNPFuncName::DPNP_FN_RECIP][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_recip_c<double>};
     fmap[DPNPFuncName::DPNP_FN_RECIP][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_recip_c<float>};
     fmap[DPNPFuncName::DPNP_FN_RECIP][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_recip_c<int>};

dpnp/dpnp_algo/dpnp_algo.pxd

@@ -63,6 +63,7 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName":  # need this na
         DPNP_FN_DOT
         DPNP_FN_EIG
         DPNP_FN_EIGVALS
+        DPNP_FN_ERF
         DPNP_FN_EXP
         DPNP_FN_EXP2
         DPNP_FN_EXPM1

dpnp/dpnp_algo/dpnp_algo.pyx

@@ -63,6 +63,7 @@ include "dpnp_algo_manipulation.pyx"
 include "dpnp_algo_mathematical.pyx"
 include "dpnp_algo_searching.pyx"
 include "dpnp_algo_sorting.pyx"
+include "dpnp_algo_special.pyx"
 include "dpnp_algo_statistics.pyx"
 include "dpnp_algo_trigonometric.pyx"
 

dpnp/dpnp_algo/dpnp_algo_special.pyx

@@ -0,0 +1,45 @@
+# cython: language_level=3
+# -*- coding: utf-8 -*-
+# *****************************************************************************
+# Copyright (c) 2016-2020, Intel Corporation
+# All rights reserved.
+#
+# Redistribution and use in source and binary forms, with or without
+# modification, are permitted provided that the following conditions are met:
+# - Redistributions of source code must retain the above copyright notice,
+#   this list of conditions and the following disclaimer.
+# - Redistributions in binary form must reproduce the above copyright notice,
+#   this list of conditions and the following disclaimer in the documentation
+#   and/or other materials provided with the distribution.
+#
+# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+# ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
+# LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+# CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+# SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+# INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+# CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+# ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
+# THE POSSIBILITY OF SUCH DAMAGE.
+# *****************************************************************************
+
+"""Module Backend (Special part)
+
+This module contains interface functions between C backend layer
+and the rest of the library
+
+"""
+
+
+from dpnp.dpnp_utils cimport *
+
+
+__all__ += [
+    'dpnp_erf',
+]
+
+
+cpdef dparray dpnp_erf(dparray x1):
+    return call_fptr_1in_1out(DPNP_FN_ERF, x1, x1.shape)

dpnp/dpnp_iface_libmath.py

@@ -39,10 +39,11 @@
 
 """
 
-
 import math
 
+from dpnp.dpnp_algo import *
 from dpnp.dparray import dparray
+from dpnp.dpnp_utils import *
 
 
 __all__ = [
@@ -76,12 +77,14 @@ def erf(in_array1):
     [0.99532227, 0.99853728, 0.99959305, 0.99989938, 0.99997791]
 
     """
+    if not use_origin_backend(in_array1):
+        if not isinstance(in_array1, dparray):
+            pass
+        else:
+            return dpnp_erf(in_array1)
 
-    if isinstance(in_array1, dparray):
-        result = dparray(in_array1.shape, dtype=in_array1.dtype)
-        for i in range(result.size):
-            result[i] = math.erf(in_array1[i])
-
-        return result
+    result = dparray(in_array1.shape, dtype=in_array1.dtype)
+    for i in range(result.size):
+        result[i] = math.erf(in_array1[i])
 
-    return math.erf(in_array1)
+    return result

tests/test_special.py

@@ -0,0 +1,29 @@
+import math
+import dpnp
+import numpy
+
+
+def test_erf():
+    a = numpy.linspace(2.0, 3.0, num=10)
+    ia = dpnp.linspace(2.0, 3.0, num=10)
+
+    numpy.testing.assert_array_equal(a, ia)
+
+    expected = numpy.empty_like(a)
+    for idx, val in enumerate(a):
+        expected[idx] = math.erf(val)
+
+    result = dpnp.erf(ia)
+
+    numpy.testing.assert_array_equal(result, expected)
+
+def test_erf_fallback():
+    a = numpy.linspace(2.0, 3.0, num=10)
+
+    expected = numpy.empty_like(a)
+    for idx, val in enumerate(a):
+        expected[idx] = math.erf(val)
+
+    result = dpnp.erf(a)
+
+    numpy.testing.assert_array_equal(result, expected)

tests_external/skipped_tests_numpy_aborted.tbl

@@ -556,6 +556,9 @@ tests/test_pocketfft.py::TestFFT1D::test_axes[fftn]
 tests/test_pocketfft.py::TestFFT1D::test_axes[ifftn]
 tests/test_pocketfft.py::TestFFT1D::test_axes[rfftn]
 tests/test_pocketfft.py::TestFFT1D::test_all_1d_norm_preserving
+tests/test_pocketfft.py::TestFFTThreadSafe::test_fft
+tests/test_pocketfft.py::TestFFTThreadSafe::test_ifft
+tests/test_pocketfft.py::TestFFTThreadSafe::test_rfft
 tests/test_randomstate_regression.py::TestRegression::test_logseries_convergence
 tests/test_regression.py::TestRegression::test_0d_string_scalar
 tests/test_regression.py::TestRegression::test_add_identity