Merge 35ddd91 into 80af9c3

Author: antonwolfy

CHANGELOG.md

@@ -12,6 +12,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 implementation of `dpnp.special.erfcx` [#2596](https://github.com/IntelPython/dpnp/pull/2596)
 
 ### 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