Added implementation of dpnp.special.erfc

CHANGELOG.md

@@ -18,6 +18,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 * 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.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)
 
 ### Changed
 

dpnp/backend/extensions/blas/CMakeLists.txt

@@ -60,8 +60,7 @@ endif()
 
 set_target_properties(${python_module_name} PROPERTIES CMAKE_POSITION_INDEPENDENT_CODE ON)
 
-target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../include)
-target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../src)
+target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../)
 target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../common)
 
 target_include_directories(${python_module_name} PUBLIC ${Dpctl_INCLUDE_DIRS})

dpnp/backend/extensions/blas/blas_py.cpp

@@ -30,6 +30,9 @@
 #include <pybind11/pybind11.h>
 #include <pybind11/stl.h>
 
+// utils extension header
+#include "ext/common.hpp"
+
 #include "dot.hpp"
 #include "dot_common.hpp"
 #include "dotc.hpp"
@@ -41,7 +44,9 @@
 namespace blas_ns = dpnp::extensions::blas;
 namespace py = pybind11;
 namespace dot_ns = blas_ns::dot;
+
 using dot_ns::dot_impl_fn_ptr_t;
+using ext::common::init_dispatch_vector;
 
 // populate dispatch vectors and tables
 void init_dispatch_vectors_tables(void)
@@ -64,7 +69,7 @@ PYBIND11_MODULE(_blas_impl, m)
     using event_vecT = std::vector<sycl::event>;
 
     {
-        dot_ns::init_dot_dispatch_vector<blas_ns::DotContigFactory>(
+        init_dispatch_vector<dot_impl_fn_ptr_t, blas_ns::DotContigFactory>(
             dot_dispatch_vector);
 
         auto dot_pyapi = [&](sycl::queue &exec_q, const arrayT &src1,
@@ -82,7 +87,7 @@ PYBIND11_MODULE(_blas_impl, m)
     }
 
     {
-        dot_ns::init_dot_dispatch_vector<blas_ns::DotcContigFactory>(
+        init_dispatch_vector<dot_impl_fn_ptr_t, blas_ns::DotcContigFactory>(
             dotc_dispatch_vector);
 
         auto dotc_pyapi = [&](sycl::queue &exec_q, const arrayT &src1,
@@ -101,7 +106,7 @@ PYBIND11_MODULE(_blas_impl, m)
     }
 
     {
-        dot_ns::init_dot_dispatch_vector<blas_ns::DotuContigFactory>(
+        init_dispatch_vector<dot_impl_fn_ptr_t, blas_ns::DotuContigFactory>(
             dotu_dispatch_vector);
 
         auto dotu_pyapi = [&](sycl::queue &exec_q, const arrayT &src1,

dpnp/backend/extensions/blas/dot_common.hpp

@@ -165,13 +165,4 @@ std::pair<sycl::event, sycl::event>
 
     return std::make_pair(args_ev, dot_ev);
 }
-
-template <template <typename fnT, typename T> typename factoryT>
-void init_dot_dispatch_vector(dot_impl_fn_ptr_t dot_dispatch_vector[])
-{
-    dpctl_td_ns::DispatchVectorBuilder<dot_impl_fn_ptr_t, factoryT,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(dot_dispatch_vector);
-}
 } // namespace dpnp::extensions::blas::dot

dpnp/backend/extensions/blas/gemm.cpp

@@ -27,6 +27,9 @@
 
 #include <pybind11/pybind11.h>
 
+// utils extension header
+#include "ext/common.hpp"
+
 // dpctl tensor headers
 #include "utils/memory_overlap.hpp"
 #include "utils/output_validation.hpp"
@@ -35,14 +38,14 @@
 #include "gemm.hpp"
 #include "types_matrix.hpp"
 
-#include "dpnp_utils.hpp"
-
 namespace dpnp::extensions::blas
 {
 namespace mkl_blas = oneapi::mkl::blas;
 namespace py = pybind11;
 namespace type_utils = dpctl::tensor::type_utils;
 
+using ext::common::init_dispatch_table;
+
 typedef sycl::event (*gemm_impl_fn_ptr_t)(sycl::queue &,
                                           oneapi::mkl::transpose,
                                           oneapi::mkl::transpose,
@@ -339,9 +342,7 @@ struct GemmContigFactory
 
 void init_gemm_dispatch_table(void)
 {
-    dpctl_td_ns::DispatchTableBuilder<gemm_impl_fn_ptr_t, GemmContigFactory,
-                                      dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_table(gemm_dispatch_table);
+    init_dispatch_table<gemm_impl_fn_ptr_t, GemmContigFactory>(
+        gemm_dispatch_table);
 }
 } // namespace dpnp::extensions::blas

dpnp/backend/extensions/blas/gemm_batch.cpp

@@ -27,6 +27,9 @@
 
 #include <pybind11/pybind11.h>
 
+// utils extension header
+#include "ext/common.hpp"
+
 // dpctl tensor headers
 #include "utils/memory_overlap.hpp"
 #include "utils/output_validation.hpp"
@@ -35,14 +38,14 @@
 #include "gemm.hpp"
 #include "types_matrix.hpp"
 
-#include "dpnp_utils.hpp"
-
 namespace dpnp::extensions::blas
 {
 namespace mkl_blas = oneapi::mkl::blas;
 namespace py = pybind11;
 namespace type_utils = dpctl::tensor::type_utils;
 
+using ext::common::init_dispatch_table;
+
 typedef sycl::event (*gemm_batch_impl_fn_ptr_t)(
     sycl::queue &,
     const std::int64_t,
@@ -415,10 +418,7 @@ struct GemmBatchContigFactory
 
 void init_gemm_batch_dispatch_table(void)
 {
-    dpctl_td_ns::DispatchTableBuilder<gemm_batch_impl_fn_ptr_t,
-                                      GemmBatchContigFactory,
-                                      dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_table(gemm_batch_dispatch_table);
+    init_dispatch_table<gemm_batch_impl_fn_ptr_t, GemmBatchContigFactory>(
+        gemm_batch_dispatch_table);
 }
 } // namespace dpnp::extensions::blas

dpnp/backend/extensions/blas/gemv.cpp

@@ -27,6 +27,9 @@
 
 #include <pybind11/pybind11.h>
 
+// utils extension header
+#include "ext/common.hpp"
+
 // dpctl tensor headers
 #include "utils/memory_overlap.hpp"
 #include "utils/output_validation.hpp"
@@ -35,14 +38,14 @@
 #include "gemv.hpp"
 #include "types_matrix.hpp"
 
-#include "dpnp_utils.hpp"
-
 namespace dpnp::extensions::blas
 {
 namespace mkl_blas = oneapi::mkl::blas;
 namespace py = pybind11;
 namespace type_utils = dpctl::tensor::type_utils;
 
+using ext::common::init_dispatch_vector;
+
 typedef sycl::event (*gemv_impl_fn_ptr_t)(sycl::queue &,
                                           oneapi::mkl::transpose,
                                           const std::int64_t,
@@ -320,9 +323,7 @@ struct GemvContigFactory
 
 void init_gemv_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<gemv_impl_fn_ptr_t, GemvContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(gemv_dispatch_vector);
+    init_dispatch_vector<gemv_impl_fn_ptr_t, GemvContigFactory>(
+        gemv_dispatch_vector);
 }
 } // namespace dpnp::extensions::blas

dpnp/backend/extensions/blas/syrk.cpp

@@ -28,6 +28,7 @@
 
 #include <pybind11/pybind11.h>
 
+// utils extension header
 #include "ext/common.hpp"
 
 // dpctl tensor headers
@@ -38,8 +39,6 @@
 #include "syrk.hpp"
 #include "types_matrix.hpp"
 
-#include "dpnp_utils.hpp"
-
 using ext::common::Align;
 
 namespace dpnp::extensions::blas
@@ -48,6 +47,8 @@ namespace mkl_blas = oneapi::mkl::blas;
 namespace py = pybind11;
 namespace type_utils = dpctl::tensor::type_utils;
 
+using ext::common::init_dispatch_vector;
+
 typedef sycl::event (*syrk_impl_fn_ptr_t)(sycl::queue &,
                                           const oneapi::mkl::transpose,
                                           const std::int64_t,
@@ -349,9 +350,7 @@ struct SyrkContigFactory
 
 void init_syrk_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<syrk_impl_fn_ptr_t, SyrkContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(syrk_dispatch_vector);
+    init_dispatch_vector<syrk_impl_fn_ptr_t, SyrkContigFactory>(
+        syrk_dispatch_vector);
 }
 } // namespace dpnp::extensions::blas

dpnp/backend/extensions/common/ext/common.hpp

@@ -30,10 +30,13 @@
 #include <pybind11/pybind11.h>
 #include <sycl/sycl.hpp>
 
+// dpctl tensor headers
 #include "utils/math_utils.hpp"
+#include "utils/type_dispatch.hpp"
 #include "utils/type_utils.hpp"
 
 namespace type_utils = dpctl::tensor::type_utils;
+namespace type_dispatch = dpctl::tensor::type_dispatch;
 
 namespace ext::common
 {
@@ -206,6 +209,25 @@ sycl::nd_range<1>
 // headers of dpctl.
 pybind11::dtype dtype_from_typenum(int dst_typenum);
 
+template <typename dispatchT,
+          template <typename fnT, typename T>
+          typename factoryT,
+          int _num_types = type_dispatch::num_types>
+inline void init_dispatch_vector(dispatchT dispatch_vector[])
+{
+    type_dispatch::DispatchVectorBuilder<dispatchT, factoryT, _num_types> dvb;
+    dvb.populate_dispatch_vector(dispatch_vector);
+}
+
+template <typename dispatchT,
+          template <typename fnT, typename D, typename S>
+          typename factoryT,
+          int _num_types = type_dispatch::num_types>
+inline void init_dispatch_table(dispatchT dispatch_table[][_num_types])
+{
+    type_dispatch::DispatchTableBuilder<dispatchT, factoryT, _num_types> dtb;
+    dtb.populate_dispatch_table(dispatch_table);
+}
 } // namespace ext::common
 
 #include "ext/details/common_internal.hpp"

dpnp/backend/extensions/indexing/CMakeLists.txt

@@ -57,8 +57,8 @@ endif()
 
 set_target_properties(${python_module_name} PROPERTIES CMAKE_POSITION_INDEPENDENT_CODE ON)
 
-target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../include)
-target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../src)
+target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../)
+target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../common)
 
 target_include_directories(${python_module_name} PUBLIC ${Dpctl_INCLUDE_DIR})
 target_include_directories(${python_module_name} PUBLIC ${Dpctl_TENSOR_INCLUDE_DIR})

dpnp/backend/extensions/indexing/choose.cpp

@@ -36,6 +36,11 @@
 
 #include "choose_kernel.hpp"
 #include "dpctl4pybind11.hpp"
+
+// utils extension header
+#include "ext/common.hpp"
+
+// dpctl tensor headers
 #include "utils/indexing_utils.hpp"
 #include "utils/memory_overlap.hpp"
 #include "utils/output_validation.hpp"
@@ -432,18 +437,13 @@ using ChooseClipFactory = ChooseFactory<fnT, IndT, T, ClipIndex<IndT>>;
 
 void init_choose_dispatch_tables(void)
 {
-    using namespace td_ns;
+    using ext::common::init_dispatch_table;
     using kernels::choose_fn_ptr_t;
 
-    DispatchTableBuilder<choose_fn_ptr_t, ChooseClipFactory, num_types>
-        dtb_choose_clip;
-    dtb_choose_clip.populate_dispatch_table(choose_clip_dispatch_table);
-
-    DispatchTableBuilder<choose_fn_ptr_t, ChooseWrapFactory, num_types>
-        dtb_choose_wrap;
-    dtb_choose_wrap.populate_dispatch_table(choose_wrap_dispatch_table);
-
-    return;
+    init_dispatch_table<choose_fn_ptr_t, ChooseClipFactory>(
+        choose_clip_dispatch_table);
+    init_dispatch_table<choose_fn_ptr_t, ChooseWrapFactory>(
+        choose_wrap_dispatch_table);
 }
 
 void init_choose(py::module_ m)

dpnp/backend/extensions/lapack/CMakeLists.txt

@@ -76,8 +76,8 @@ endif()
 
 set_target_properties(${python_module_name} PROPERTIES CMAKE_POSITION_INDEPENDENT_CODE ON)
 
-target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../include)
-target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../src)
+target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../)
+target_include_directories(${python_module_name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../common)
 
 target_include_directories(${python_module_name} PUBLIC ${Dpctl_INCLUDE_DIR})
 target_include_directories(${python_module_name} PUBLIC ${Dpctl_TENSOR_INCLUDE_DIR})

dpnp/backend/extensions/lapack/evd_common_utils.hpp

@@ -30,25 +30,11 @@
 // dpctl tensor headers
 #include "utils/memory_overlap.hpp"
 #include "utils/output_validation.hpp"
-#include "utils/type_dispatch.hpp"
 
 namespace dpnp::extensions::lapack::evd
 {
-namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
 namespace py = pybind11;
 
-template <typename dispatchT,
-          template <typename fnT, typename T, typename RealT>
-          typename factoryT>
-void init_evd_dispatch_table(
-    dispatchT evd_dispatch_table[][dpctl_td_ns::num_types])
-{
-    dpctl_td_ns::DispatchTableBuilder<dispatchT, factoryT,
-                                      dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_table(evd_dispatch_table);
-}
-
 inline void common_evd_checks(sycl::queue &exec_q,
                               const dpctl::tensor::usm_ndarray &eig_vecs,
                               const dpctl::tensor::usm_ndarray &eig_vals,

dpnp/backend/extensions/lapack/geqrf.cpp

@@ -27,6 +27,9 @@
 
 #include <pybind11/pybind11.h>
 
+// utils extension header
+#include "ext/common.hpp"
+
 // dpctl tensor headers
 #include "utils/memory_overlap.hpp"
 #include "utils/sycl_alloc_utils.hpp"
@@ -35,14 +38,14 @@
 #include "geqrf.hpp"
 #include "types_matrix.hpp"
 
-#include "dpnp_utils.hpp"
-
 namespace dpnp::extensions::lapack
 {
 namespace mkl_lapack = oneapi::mkl::lapack;
 namespace py = pybind11;
 namespace type_utils = dpctl::tensor::type_utils;
 
+using ext::common::init_dispatch_vector;
+
 typedef sycl::event (*geqrf_impl_fn_ptr_t)(sycl::queue &,
                                            const std::int64_t,
                                            const std::int64_t,
@@ -250,9 +253,7 @@ struct GeqrfContigFactory
 
 void init_geqrf_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<geqrf_impl_fn_ptr_t, GeqrfContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(geqrf_dispatch_vector);
+    init_dispatch_vector<geqrf_impl_fn_ptr_t, GeqrfContigFactory>(
+        geqrf_dispatch_vector);
 }
 } // namespace dpnp::extensions::lapack