Add implementation of dpnp.scipy.special.erfcx (#2596)

The PR adds the implementation of `dpnp.special.erfcx` function,
including tests coverage.

It also includes the special implementation of `erfcx` which will be
used when building dpnp from the source with support of AMD or NVIDIA
devices.

Author: antonwolfy

CHANGELOG.md

@@ -13,6 +13,7 @@ This release changes the license from `BSD-2-Clause` to `BSD-3-Clause`.
 * Added the docstrings to `dpnp.linalg.LinAlgError` exception [#2613](https://github.com/IntelPython/dpnp/pull/2613)
 * Added implementation of `dpnp.linalg.lu_solve` for batch inputs (SciPy-compatible) [#2619](https://github.com/IntelPython/dpnp/pull/2619)
 * Added `dpnp.exceptions` submodule to aggregate the generic exceptions used by dpnp [#2616](https://github.com/IntelPython/dpnp/pull/2616)
+* Added implementation of `dpnp.scipy.special.erfcx` [#2596](https://github.com/IntelPython/dpnp/pull/2596)
 
 ### Changed
 
@@ -44,10 +45,10 @@ This release is compatible with NumPy 2.3.3.
 * Added `timeout-minutes` property to GitHub jobs [#2526](https://github.com/IntelPython/dpnp/pull/2526)
 * Added implementation of `dpnp.ndarray.data` and `dpnp.ndarray.data.ptr` attributes [#2521](https://github.com/IntelPython/dpnp/pull/2521)
 * Added `dpnp.ndarray.__contains__` method [#2534](https://github.com/IntelPython/dpnp/pull/2534)
-* Added implementation of `dpnp.linalg.lu_factor` (SciPy-compatible) [#2557](https://github.com/IntelPython/dpnp/pull/2557), [#2565](https://github.com/IntelPython/dpnp/pull/2565)
+* Added implementation of `dpnp.scipy.linalg.lu_factor` (SciPy-compatible) [#2557](https://github.com/IntelPython/dpnp/pull/2557), [#2565](https://github.com/IntelPython/dpnp/pull/2565)
 * Added implementation of `dpnp.piecewise` [#2550](https://github.com/IntelPython/dpnp/pull/2550)
-* Added implementation of `dpnp.linalg.lu_solve` for 2D inputs (SciPy-compatible) [#2575](https://github.com/IntelPython/dpnp/pull/2575)
-* Added implementation of `dpnp.special.erfc` [#2588](https://github.com/IntelPython/dpnp/pull/2588)
+* Added implementation of `dpnp.scipy.linalg.lu_solve` for 2D inputs (SciPy-compatible) [#2575](https://github.com/IntelPython/dpnp/pull/2575)
+* Added implementation of `dpnp.scipy.special.erfc` [#2588](https://github.com/IntelPython/dpnp/pull/2588)
 * Added `dpnp.scipy` submodule to aggregate new SciPy-compatible functions from `linalg` and `special` namespaces [#2603](https://github.com/IntelPython/dpnp/pull/2603)
 
 ### Changed

doc/reference/scipy_special.rst

@@ -15,6 +15,5 @@ Error function and Fresnel integrals
    erf
    erfc
    erfcx
-   erfi
    erfinv
    erfcinv

dpnp/backend/extensions/ufunc/elementwise_functions/erf_funcs.cpp

@@ -212,6 +212,7 @@ static void populate(py::module_ m,
 
 MACRO_DEFINE_IMPL(erf, Erf);
 MACRO_DEFINE_IMPL(erfc, Erfc);
+MACRO_DEFINE_IMPL(erfcx, Erfcx);
 } // namespace impl
 
 void init_erf_funcs(py::module_ m)
@@ -231,5 +232,9 @@ void init_erf_funcs(py::module_ m)
     impl::populate<impl::ErfcContigFactory, impl::ErfcStridedFactory>(
         m, "_erfc", "", impl::erfc_contig_dispatch_vector,
         impl::erfc_strided_dispatch_vector);
+
+    impl::populate<impl::ErfcxContigFactory, impl::ErfcxStridedFactory>(
+        m, "_erfcx", "", impl::erfcx_contig_dispatch_vector,
+        impl::erfcx_strided_dispatch_vector);
 }
 } // namespace dpnp::extensions::ufunc

dpnp/backend/extensions/vm/erf_funcs.cpp

@@ -133,6 +133,7 @@ using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
 
 MACRO_DEFINE_IMPL(erf, Erf);
 MACRO_DEFINE_IMPL(erfc, Erfc);
+MACRO_DEFINE_IMPL(erfcx, Erfcx);
 
 template <template <typename fnT, typename T> typename factoryT>
 static void populate(py::module_ m,
@@ -187,5 +188,11 @@ void init_erf_funcs(py::module_ m)
         "Call `erfc` function from OneMKL VM library to compute the "
         "complementary error function value of vector elements",
         impl::erfc_contig_dispatch_vector);
+
+    impl::populate<impl::ErfcxContigFactory>(
+        m, "_erfcx",
+        "Call `erfcx` function from OneMKL VM library to compute the scaled "
+        "complementary error function value of vector elements",
+        impl::erfcx_contig_dispatch_vector);
 }
 } // namespace dpnp::extensions::vm

dpnp/backend/kernels/elementwise_functions/erf.hpp

@@ -32,6 +32,19 @@
 
 #include <sycl/sycl.hpp>
 
+/**
+ * Include <sycl/ext/intel/math.hpp> only when targeting to Intel devices.
+ */
+#if (defined(__SPIR__) || defined(__SPIRV__)) && defined(__INTEL_LLVM_COMPILER)
+#define __SYCL_EXT_INTEL_MATH_SUPPORT
+#endif
+
+#if defined(__SYCL_EXT_INTEL_MATH_SUPPORT)
+#include <sycl/ext/intel/math.hpp>
+#else
+#include "erfcx.hpp"
+#endif
+
 namespace dpnp::kernels::erfs
 {
 template <typename OpT, typename ArgT, typename ResT>
@@ -65,13 +78,20 @@ struct BaseFunctor
         template <typename Tp>                                                 \
         static Tp apply(const Tp &x)                                           \
         {                                                                      \
-            return sycl::__name__(x);                                          \
+            return __name__(x);                                                \
         }                                                                      \
     };                                                                         \
                                                                                \
     template <typename ArgT, typename ResT>                                    \
     using __f_name__##Functor = BaseFunctor<__f_name__##Op, ArgT, ResT>;
 
-MACRO_DEFINE_FUNCTOR(erf, Erf);
-MACRO_DEFINE_FUNCTOR(erfc, Erfc);
+MACRO_DEFINE_FUNCTOR(sycl::erf, Erf);
+MACRO_DEFINE_FUNCTOR(sycl::erfc, Erfc);
+MACRO_DEFINE_FUNCTOR(
+#if defined(__SYCL_EXT_INTEL_MATH_SUPPORT)
+    sycl::ext::intel::math::erfcx,
+#else
+    impl::erfcx,
+#endif
+    Erfcx);
 } // namespace dpnp::kernels::erfs