Implement dpnp.erf through pybind11 extension

.github/workflows/check-onemath.yaml

@@ -15,6 +15,7 @@ env:
   build-with-oneapi-env: 'environments/build_with_oneapi.yml'
   dpctl-pkg-env: 'environments/dpctl_pkg.yml'
   oneapi-pkgs-env: 'environments/oneapi_pkgs.yml'
+  test-pkg-env: 'environments/test.yml'
   test-env-name: 'test_onemath'
   rerun-tests-on-failure: 'true'
   rerun-tests-max-attempts: 2
@@ -49,7 +50,10 @@ jobs:
 
       - name: Merge conda env files
         run: |
-          conda-merge ${{ env.dpctl-pkg-env }} ${{ env.oneapi-pkgs-env }} ${{ env.build-with-oneapi-env }} > ${{ env.environment-file }}
+          conda-merge ${{ env.dpctl-pkg-env }}         \
+                      ${{ env.oneapi-pkgs-env }}       \
+                      ${{ env.build-with-oneapi-env }} \
+                      ${{ env.test-pkg-env }} > ${{ env.environment-file }}
           cat ${{ env.environment-file }}
 
       - name: Upload artifact

.github/workflows/cron-run-tests.yaml

@@ -84,12 +84,12 @@ jobs:
         id: install_dpnp
         continue-on-error: true
         run: |
-          mamba install ${{ env.package-name }}=${{ steps.find_latest_tag.outputs.tag }} pytest ${{ env.channels-list }}
+          mamba install ${{ env.package-name }}=${{ steps.find_latest_tag.outputs.tag }} pytest scipy ${{ env.channels-list }}
 
       - name: ReInstall dpnp
         if: steps.install_dpnp.outcome == 'failure'
         run: |
-          mamba install ${{ env.package-name }}=${{ steps.find_latest_tag.outputs.tag }} pytest ${{ env.channels-list }}
+          mamba install ${{ env.package-name }}=${{ steps.find_latest_tag.outputs.tag }} pytest scipy ${{ env.channels-list }}
 
       - name: List installed packages
         run: mamba list

.github/workflows/generate_coverage.yaml

@@ -25,6 +25,7 @@ jobs:
       environment-file: 'environments/environment.yml'
       build-with-oneapi-env: 'environments/build_with_oneapi.yml'
       coverage-env: 'environments/coverage.yml'
+      test-pkg-env: 'environments/test.yml'
       oneapi-pkgs-env: ''
       # Enable env when it's required to use only conda packages without OneAPI installation
       # oneapi-pkgs-env: '${{ github.workspace }}/environments/oneapi_pkgs.yml'
@@ -73,7 +74,10 @@ jobs:
 
       - name: Merge conda env files
         run: |
-          conda-merge ${{ env.build-with-oneapi-env }} ${{ env.coverage-env }} ${{ env.oneapi-pkgs-env }} > ${{ env.environment-file }}
+          conda-merge ${{ env.build-with-oneapi-env }} \
+                      ${{ env.coverage-env }}          \
+                      ${{ env.oneapi-pkgs-env }}       \
+                      ${{ env.test-pkg-env }} > ${{ env.environment-file }}
           cat ${{ env.environment-file }}
 
       - name: Setup miniconda

.pre-commit-config.yaml

@@ -113,7 +113,7 @@ repos:
             "--disable=redefined-builtin",
             "--disable=unused-wildcard-import"
             ]
-        files: '^dpnp/(dpnp_iface.*|fft|linalg|dpnp_array)'
+        files: '^dpnp/(dpnp_iface.*|fft|linalg|special|dpnp_array)'
 -   repo: https://github.com/macisamuele/language-formatters-pre-commit-hooks
     rev: v2.15.0
     hooks:

CHANGELOG.md

@@ -29,7 +29,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 * Aligned the license expression with `PEP-639` [#2511](https://github.com/IntelPython/dpnp/pull/2511)
 * Bumped oneMKL version up to `v0.8` [#2514](https://github.com/IntelPython/dpnp/pull/2514)
 * Removed the use of class template argument deduction for alias template to conform to the C++17 standard [#2517](https://github.com/IntelPython/dpnp/pull/2517)
-* Changed th order of individual FFTs over `axes` for `dpnp.fft.irfftn` to be in forward order [#2524](https://github.com/IntelPython/dpnp/pull/2524)
+* Changed the order of individual FFTs over `axes` for `dpnp.fft.irfftn` to be in forward order [#2524](https://github.com/IntelPython/dpnp/pull/2524)
 * Replaced the use of `numpy.testing.suppress_warnings` with appropriate calls from the warnings module [#2529](https://github.com/IntelPython/dpnp/pull/2529)
 * Improved documentations of `dpnp.ndarray` class and added a page with description of supported constants [#2422](https://github.com/IntelPython/dpnp/pull/2422)
 * Updated `dpnp.size` to accept tuple of ints for `axes` argument [#2536](https://github.com/IntelPython/dpnp/pull/2536)
@@ -41,6 +41,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 * Refactored backend implementation of `dpnp.linalg.solve` to use oneMKL LAPACK `gesv` directly [#2558](https://github.com/IntelPython/dpnp/pull/2558)
 * Improved performance of `dpnp.isclose` function by implementing a dedicated kernel for scalar `rtol` and `atol` arguments [#2540](https://github.com/IntelPython/dpnp/pull/2540)
 * Extended `dpnp.pad` to support `pad_width` keyword as a dictionary [#2535](https://github.com/IntelPython/dpnp/pull/2535)
+* Redesigned `dpnp.erf` function through pybind11 extension of OneMKL call or dedicated kernel in `ufunc` namespace [#2551](https://github.com/IntelPython/dpnp/pull/2551)
 
 ### Deprecated
 

conda-recipe/meta.yaml

@@ -63,6 +63,7 @@ test:
     requires:
       - pytest
       - setuptools
+      - scipy
 
 about:
     home: https://github.com/IntelPython/dpnp

doc/reference/special.rst

@@ -1,17 +1,20 @@
-Special Functions
-=================
+.. currentmodule:: dpnp.special
 
-.. https://docs.scipy.org/doc/scipy/reference/special.html
+Special functions (:mod:`dpnp.special`)
+=======================================
 
-Error Function
---------------
+.. Hint:: `SciPy API Reference: Special functions (scipy.special) <https://docs.scipy.org/doc/scipy/reference/special.html>`_
+
+Error function and Fresnel integrals
+------------------------------------
 
 .. autosummary::
    :toctree: generated/
    :nosignatures:
 
-   dpnp.erf
-   dpnp.erfc
-   dpnp.erfcx
-   dpnp.erfinv
-   dpnp.erfcinv
+   erf
+   erfc
+   erfcx
+   erfi
+   erfinv
+   erfcinv

dpnp/backend/CMakeLists.txt

@@ -26,7 +26,6 @@
 set(DPNP_SRC
     kernels/dpnp_krnl_arraycreation.cpp
     kernels/dpnp_krnl_common.cpp
-    kernels/dpnp_krnl_elemwise.cpp
     kernels/dpnp_krnl_mathematical.cpp
     kernels/dpnp_krnl_random.cpp
     kernels/dpnp_krnl_sorting.cpp

dpnp/backend/extensions/ufunc/CMakeLists.txt

@@ -27,6 +27,7 @@ set(_elementwise_sources
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/bitwise_count.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/common.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/degrees.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/erf.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fabs.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fix.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/float_power.cpp

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

@@ -27,6 +27,7 @@
 
 #include "bitwise_count.hpp"
 #include "degrees.hpp"
+#include "erf.hpp"
 #include "fabs.hpp"
 #include "fix.hpp"
 #include "float_power.hpp"
@@ -57,6 +58,7 @@ void init_elementwise_functions(py::module_ m)
 {
     init_bitwise_count(m);
     init_degrees(m);
+    init_erf(m);
     init_fabs(m);
     init_fix(m);
     init_float_power(m);

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

@@ -0,0 +1,132 @@
+//*****************************************************************************
+// Copyright (c) 2025, 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.
+//*****************************************************************************
+
+#include <type_traits>
+#include <vector>
+
+#include <sycl/sycl.hpp>
+
+#include "dpctl4pybind11.hpp"
+
+#include "erf.hpp"
+#include "kernels/elementwise_functions/erf.hpp"
+#include "populate.hpp"
+
+// include a local copy of elementwise common header from dpctl tensor:
+// dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp
+// TODO: replace by including dpctl header once available
+#include "../../elementwise_functions/elementwise_functions.hpp"
+
+// dpctl tensor headers
+#include "kernels/elementwise_functions/common.hpp"
+#include "utils/type_dispatch.hpp"
+
+namespace dpnp::extensions::ufunc
+{
+namespace py = pybind11;
+namespace py_int = dpnp::extensions::py_internal;
+
+namespace impl
+{
+namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace td_ns = dpctl::tensor::type_dispatch;
+
+/**
+ * @brief A factory to define pairs of supported types for which
+ * sycl::erf<T> function is available.
+ *
+ * @tparam T Type of input vector `a` and of result vector `y`.
+ */
+template <typename T>
+struct OutputType
+{
+    /**
+     * scipy>=1.16 assumes a pair 'e->d', but dpnp 'e->f' without an extra
+     * kernel 'e->d' (when fp64 supported) to reduce memory footprint
+     */
+    using value_type = typename std::disjunction<
+        td_ns::TypeMapResultEntry<T, sycl::half, float>,
+        td_ns::TypeMapResultEntry<T, float>,
+        td_ns::TypeMapResultEntry<T, double>,
+        td_ns::DefaultResultEntry<void>>::result_type;
+};
+
+using dpnp::kernels::erf::ErfFunctor;
+
+template <typename argT,
+          typename resT = argT,
+          unsigned int vec_sz = 4,
+          unsigned int n_vecs = 2,
+          bool enable_sg_loadstore = true>
+using ContigFunctor = ew_cmn_ns::UnaryContigFunctor<argT,
+                                                    resT,
+                                                    ErfFunctor<argT, resT>,
+                                                    vec_sz,
+                                                    n_vecs,
+                                                    enable_sg_loadstore>;
+
+template <typename argTy, typename resTy, typename IndexerT>
+using StridedFunctor = ew_cmn_ns::
+    UnaryStridedFunctor<argTy, resTy, IndexerT, ErfFunctor<argTy, resTy>>;
+
+using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
+using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
+
+static unary_contig_impl_fn_ptr_t erf_contig_dispatch_vector[td_ns::num_types];
+static int erf_output_typeid_vector[td_ns::num_types];
+static unary_strided_impl_fn_ptr_t
+    erf_strided_dispatch_vector[td_ns::num_types];
+
+MACRO_POPULATE_DISPATCH_VECTORS(erf);
+} // namespace impl
+
+void init_erf(py::module_ m)
+{
+    using arrayT = dpctl::tensor::usm_ndarray;
+    using event_vecT = std::vector<sycl::event>;
+    {
+        impl::populate_erf_dispatch_vectors();
+        using impl::erf_contig_dispatch_vector;
+        using impl::erf_output_typeid_vector;
+        using impl::erf_strided_dispatch_vector;
+
+        auto erf_pyapi = [&](const arrayT &src, const arrayT &dst,
+                             sycl::queue &exec_q,
+                             const event_vecT &depends = {}) {
+            return py_int::py_unary_ufunc(
+                src, dst, exec_q, depends, erf_output_typeid_vector,
+                erf_contig_dispatch_vector, erf_strided_dispatch_vector);
+        };
+        m.def("_erf", erf_pyapi, "", py::arg("src"), py::arg("dst"),
+              py::arg("sycl_queue"), py::arg("depends") = py::list());
+
+        auto erf_result_type_pyapi = [&](const py::dtype &dtype) {
+            return py_int::py_unary_ufunc_result_type(dtype,
+                                                      erf_output_typeid_vector);
+        };
+        m.def("_erf_result_type", erf_result_type_pyapi);
+    }
+}
+} // namespace dpnp::extensions::ufunc

dpnp/backend/extensions/ufunc/elementwise_functions/erf.hpp

@@ -0,0 +1,35 @@
+//*****************************************************************************
+// Copyright (c) 2025, 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.
+//*****************************************************************************
+
+#pragma once
+
+#include <pybind11/pybind11.h>
+
+namespace py = pybind11;
+
+namespace dpnp::extensions::ufunc
+{
+void init_erf(py::module_ m);
+} // namespace dpnp::extensions::ufunc

dpnp/backend/extensions/vm/CMakeLists.txt

@@ -42,6 +42,7 @@ if(NOT _use_onemath)
         ${CMAKE_CURRENT_SOURCE_DIR}/cos.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/cosh.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/div.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/erf.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/exp.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/exp2.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/expm1.cpp