diff --git a/CHANGELOG.md b/CHANGELOG.md
index ff54b2e6d151..376ed6497a38 100644
--- a/CHANGELOG.md
+++ b/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
 
diff --git a/dpnp/backend/extensions/blas/CMakeLists.txt b/dpnp/backend/extensions/blas/CMakeLists.txt
index 3f81253169d1..f5e0ecc737de 100644
--- a/dpnp/backend/extensions/blas/CMakeLists.txt
+++ b/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})
diff --git a/dpnp/backend/extensions/blas/blas_py.cpp b/dpnp/backend/extensions/blas/blas_py.cpp
index 850c1e784c7b..36fc29bae482 100644
--- a/dpnp/backend/extensions/blas/blas_py.cpp
+++ b/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,
diff --git a/dpnp/backend/extensions/blas/dot_common.hpp b/dpnp/backend/extensions/blas/dot_common.hpp
index 169421a2464c..faa31cb6b388 100644
--- a/dpnp/backend/extensions/blas/dot_common.hpp
+++ b/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
diff --git a/dpnp/backend/extensions/blas/gemm.cpp b/dpnp/backend/extensions/blas/gemm.cpp
index 5a411d94ba1f..4d3a98afec66 100644
--- a/dpnp/backend/extensions/blas/gemm.cpp
+++ b/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
diff --git a/dpnp/backend/extensions/blas/gemm_batch.cpp b/dpnp/backend/extensions/blas/gemm_batch.cpp
index f329f8f4be58..093b7d630b9d 100644
--- a/dpnp/backend/extensions/blas/gemm_batch.cpp
+++ b/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
diff --git a/dpnp/backend/extensions/blas/gemv.cpp b/dpnp/backend/extensions/blas/gemv.cpp
index 37da858c8e69..3e535cee2c36 100644
--- a/dpnp/backend/extensions/blas/gemv.cpp
+++ b/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
diff --git a/dpnp/backend/extensions/blas/syrk.cpp b/dpnp/backend/extensions/blas/syrk.cpp
index d432a810ce6b..f016b90a73ab 100644
--- a/dpnp/backend/extensions/blas/syrk.cpp
+++ b/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
diff --git a/dpnp/backend/extensions/common/ext/common.hpp b/dpnp/backend/extensions/common/ext/common.hpp
index 080df62b25e3..1c1216670425 100644
--- a/dpnp/backend/extensions/common/ext/common.hpp
+++ b/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"
diff --git a/dpnp/backend/extensions/indexing/CMakeLists.txt b/dpnp/backend/extensions/indexing/CMakeLists.txt
index 527097d4ffdf..d9de70ec1010 100644
--- a/dpnp/backend/extensions/indexing/CMakeLists.txt
+++ b/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})
diff --git a/dpnp/backend/extensions/indexing/choose.cpp b/dpnp/backend/extensions/indexing/choose.cpp
index e0fd0767aca2..78794bab44b5 100644
--- a/dpnp/backend/extensions/indexing/choose.cpp
+++ b/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)
diff --git a/dpnp/backend/extensions/lapack/CMakeLists.txt b/dpnp/backend/extensions/lapack/CMakeLists.txt
index 1f35ee4d47c1..fb5e55faa01d 100644
--- a/dpnp/backend/extensions/lapack/CMakeLists.txt
+++ b/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})
diff --git a/dpnp/backend/extensions/lapack/evd_common_utils.hpp b/dpnp/backend/extensions/lapack/evd_common_utils.hpp
index 8525b275f436..3701438c98fd 100644
--- a/dpnp/backend/extensions/lapack/evd_common_utils.hpp
+++ b/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,
diff --git a/dpnp/backend/extensions/lapack/geqrf.cpp b/dpnp/backend/extensions/lapack/geqrf.cpp
index 4f7c7e64a19e..03b8f25501e4 100644
--- a/dpnp/backend/extensions/lapack/geqrf.cpp
+++ b/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
diff --git a/dpnp/backend/extensions/lapack/geqrf_batch.cpp b/dpnp/backend/extensions/lapack/geqrf_batch.cpp
index 9783b9f55df2..98104cf65128 100644
--- a/dpnp/backend/extensions/lapack/geqrf_batch.cpp
+++ b/dpnp/backend/extensions/lapack/geqrf_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/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_batch_impl_fn_ptr_t)(
     sycl::queue &,
     std::int64_t,
@@ -260,10 +263,7 @@ struct GeqrfBatchContigFactory
 
 void init_geqrf_batch_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<geqrf_batch_impl_fn_ptr_t,
-                                       GeqrfBatchContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(geqrf_batch_dispatch_vector);
+    init_dispatch_vector<geqrf_batch_impl_fn_ptr_t, GeqrfBatchContigFactory>(
+        geqrf_batch_dispatch_vector);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/lapack/gesv.cpp b/dpnp/backend/extensions/lapack/gesv.cpp
index 7697316d02e9..22154be9b3a5 100644
--- a/dpnp/backend/extensions/lapack/gesv.cpp
+++ b/dpnp/backend/extensions/lapack/gesv.cpp
@@ -28,6 +28,9 @@
 
 #include <pybind11/pybind11.h>
 
+// utils extension header
+#include "ext/common.hpp"
+
 // dpctl tensor headers
 #include "utils/type_utils.hpp"
 
@@ -43,6 +46,7 @@ namespace py = pybind11;
 namespace type_utils = dpctl::tensor::type_utils;
 
 using dpctl::tensor::alloc_utils::sycl_free_noexcept;
+using ext::common::init_dispatch_vector;
 
 typedef sycl::event (*gesv_impl_fn_ptr_t)(sycl::queue &,
                                           const std::int64_t,
@@ -284,9 +288,7 @@ struct GesvContigFactory
 
 void init_gesv_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<gesv_impl_fn_ptr_t, GesvContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(gesv_dispatch_vector);
+    init_dispatch_vector<gesv_impl_fn_ptr_t, GesvContigFactory>(
+        gesv_dispatch_vector);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/lapack/gesv_batch.cpp b/dpnp/backend/extensions/lapack/gesv_batch.cpp
index dd9368d7132d..d3ca2510efbe 100644
--- a/dpnp/backend/extensions/lapack/gesv_batch.cpp
+++ b/dpnp/backend/extensions/lapack/gesv_batch.cpp
@@ -28,6 +28,9 @@
 
 #include <pybind11/pybind11.h>
 
+// utils extension header
+#include "ext/common.hpp"
+
 // dpctl tensor headers
 #include "utils/type_utils.hpp"
 
@@ -43,6 +46,7 @@ namespace py = pybind11;
 namespace type_utils = dpctl::tensor::type_utils;
 
 using dpctl::tensor::alloc_utils::sycl_free_noexcept;
+using ext::common::init_dispatch_vector;
 
 typedef sycl::event (*gesv_batch_impl_fn_ptr_t)(
     sycl::queue &,
@@ -425,10 +429,7 @@ struct GesvBatchContigFactory
 
 void init_gesv_batch_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<gesv_batch_impl_fn_ptr_t,
-                                       GesvBatchContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(gesv_batch_dispatch_vector);
+    init_dispatch_vector<gesv_batch_impl_fn_ptr_t, GesvBatchContigFactory>(
+        gesv_batch_dispatch_vector);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/lapack/gesvd.cpp b/dpnp/backend/extensions/lapack/gesvd.cpp
index 16840b2be8f8..c3b00f66ea46 100644
--- a/dpnp/backend/extensions/lapack/gesvd.cpp
+++ b/dpnp/backend/extensions/lapack/gesvd.cpp
@@ -27,6 +27,9 @@
 
 #include <pybind11/pybind11.h>
 
+// utils extension header
+#include "ext/common.hpp"
+
 // dpctl tensor headers
 #include "utils/type_utils.hpp"
 
@@ -41,6 +44,8 @@ namespace mkl_lapack = oneapi::mkl::lapack;
 namespace py = pybind11;
 namespace type_utils = dpctl::tensor::type_utils;
 
+using ext::common::init_dispatch_table;
+
 typedef sycl::event (*gesvd_impl_fn_ptr_t)(sycl::queue &,
                                            const oneapi::mkl::jobsvd,
                                            const oneapi::mkl::jobsvd,
@@ -227,9 +232,7 @@ struct GesvdContigFactory
 
 void init_gesvd_dispatch_table(void)
 {
-    dpctl_td_ns::DispatchTableBuilder<gesvd_impl_fn_ptr_t, GesvdContigFactory,
-                                      dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_table(gesvd_dispatch_table);
+    init_dispatch_table<gesvd_impl_fn_ptr_t, GesvdContigFactory>(
+        gesvd_dispatch_table);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/lapack/gesvd_batch.cpp b/dpnp/backend/extensions/lapack/gesvd_batch.cpp
index f2c9db43895b..7a031e472cd9 100644
--- a/dpnp/backend/extensions/lapack/gesvd_batch.cpp
+++ b/dpnp/backend/extensions/lapack/gesvd_batch.cpp
@@ -27,6 +27,9 @@
 
 #include <pybind11/pybind11.h>
 
+// utils extension header
+#include "ext/common.hpp"
+
 // dpctl tensor headers
 #include "utils/type_utils.hpp"
 
@@ -41,6 +44,8 @@ namespace mkl_lapack = oneapi::mkl::lapack;
 namespace py = pybind11;
 namespace type_utils = dpctl::tensor::type_utils;
 
+using ext::common::init_dispatch_table;
+
 typedef sycl::event (*gesvd_batch_impl_fn_ptr_t)(
     sycl::queue &,
     const oneapi::mkl::jobsvd,
@@ -296,10 +301,7 @@ struct GesvdBatchContigFactory
 
 void init_gesvd_batch_dispatch_table(void)
 {
-    dpctl_td_ns::DispatchTableBuilder<gesvd_batch_impl_fn_ptr_t,
-                                      GesvdBatchContigFactory,
-                                      dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_table(gesvd_batch_dispatch_table);
+    init_dispatch_table<gesvd_batch_impl_fn_ptr_t, GesvdBatchContigFactory>(
+        gesvd_batch_dispatch_table);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/lapack/getrf.cpp b/dpnp/backend/extensions/lapack/getrf.cpp
index 81231b5055af..aae089d7f7e6 100644
--- a/dpnp/backend/extensions/lapack/getrf.cpp
+++ b/dpnp/backend/extensions/lapack/getrf.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 "getrf.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 (*getrf_impl_fn_ptr_t)(sycl::queue &,
                                            const std::int64_t,
                                            const std::int64_t,
@@ -262,9 +265,7 @@ struct GetrfContigFactory
 
 void init_getrf_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<getrf_impl_fn_ptr_t, GetrfContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(getrf_dispatch_vector);
+    init_dispatch_vector<getrf_impl_fn_ptr_t, GetrfContigFactory>(
+        getrf_dispatch_vector);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/lapack/getrf_batch.cpp b/dpnp/backend/extensions/lapack/getrf_batch.cpp
index 446f565d6e49..ad5977fb6494 100644
--- a/dpnp/backend/extensions/lapack/getrf_batch.cpp
+++ b/dpnp/backend/extensions/lapack/getrf_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/sycl_alloc_utils.hpp"
@@ -35,14 +38,14 @@
 #include "getrf.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 (*getrf_batch_impl_fn_ptr_t)(
     sycl::queue &,
     std::int64_t,
@@ -295,10 +298,7 @@ struct GetrfBatchContigFactory
 
 void init_getrf_batch_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<getrf_batch_impl_fn_ptr_t,
-                                       GetrfBatchContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(getrf_batch_dispatch_vector);
+    init_dispatch_vector<getrf_batch_impl_fn_ptr_t, GetrfBatchContigFactory>(
+        getrf_batch_dispatch_vector);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/lapack/getri_batch.cpp b/dpnp/backend/extensions/lapack/getri_batch.cpp
index a59e7e3ae606..8293da2bc179 100644
--- a/dpnp/backend/extensions/lapack/getri_batch.cpp
+++ b/dpnp/backend/extensions/lapack/getri_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/sycl_alloc_utils.hpp"
@@ -35,14 +38,14 @@
 #include "getri.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 (*getri_batch_impl_fn_ptr_t)(
     sycl::queue &,
     std::int64_t,
@@ -282,10 +285,7 @@ struct GetriBatchContigFactory
 
 void init_getri_batch_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<getri_batch_impl_fn_ptr_t,
-                                       GetriBatchContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(getri_batch_dispatch_vector);
+    init_dispatch_vector<getri_batch_impl_fn_ptr_t, GetriBatchContigFactory>(
+        getri_batch_dispatch_vector);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/lapack/getrs.cpp b/dpnp/backend/extensions/lapack/getrs.cpp
index 8185f3a06a79..4112a7c95ae7 100644
--- a/dpnp/backend/extensions/lapack/getrs.cpp
+++ b/dpnp/backend/extensions/lapack/getrs.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"
@@ -36,14 +39,14 @@
 #include "linalg_exceptions.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 (*getrs_impl_fn_ptr_t)(sycl::queue &,
                                            oneapi::mkl::transpose,
                                            const std::int64_t,
@@ -298,9 +301,7 @@ struct GetrsContigFactory
 
 void init_getrs_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<getrs_impl_fn_ptr_t, GetrsContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(getrs_dispatch_vector);
+    init_dispatch_vector<getrs_impl_fn_ptr_t, GetrsContigFactory>(
+        getrs_dispatch_vector);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/lapack/heevd.cpp b/dpnp/backend/extensions/lapack/heevd.cpp
index 19d875951c9b..9f5d7617ef28 100644
--- a/dpnp/backend/extensions/lapack/heevd.cpp
+++ b/dpnp/backend/extensions/lapack/heevd.cpp
@@ -30,6 +30,9 @@
 #include "evd_common.hpp"
 #include "heevd.hpp"
 
+// utils extension header
+#include "ext/common.hpp"
+
 // dpctl tensor headers
 #include "utils/type_utils.hpp"
 
@@ -38,6 +41,8 @@ namespace dpnp::extensions::lapack
 namespace mkl_lapack = oneapi::mkl::lapack;
 namespace type_utils = dpctl::tensor::type_utils;
 
+using ext::common::init_dispatch_table;
+
 template <typename T, typename RealT>
 static sycl::event heevd_impl(sycl::queue &exec_q,
                               const oneapi::mkl::job jobz,
@@ -137,7 +142,7 @@ void init_heevd(py::module_ m)
                                                  [dpctl_td_ns::num_types];
 
     {
-        evd::init_evd_dispatch_table<evd_impl_fn_ptr_t, HeevdContigFactory>(
+        init_dispatch_table<evd_impl_fn_ptr_t, HeevdContigFactory>(
             heevd_dispatch_table);
 
         auto heevd_pyapi = [&](sycl::queue &exec_q, const std::int8_t jobz,
diff --git a/dpnp/backend/extensions/lapack/heevd_batch.cpp b/dpnp/backend/extensions/lapack/heevd_batch.cpp
index d700e7505da5..de554cb8f8e6 100644
--- a/dpnp/backend/extensions/lapack/heevd_batch.cpp
+++ b/dpnp/backend/extensions/lapack/heevd_batch.cpp
@@ -31,6 +31,9 @@
 #include "evd_batch_common.hpp"
 #include "heevd_batch.hpp"
 
+// utils extension header
+#include "ext/common.hpp"
+
 // dpctl tensor headers
 #include "utils/type_utils.hpp"
 
@@ -39,6 +42,8 @@ namespace dpnp::extensions::lapack
 namespace mkl_lapack = oneapi::mkl::lapack;
 namespace type_utils = dpctl::tensor::type_utils;
 
+using ext::common::init_dispatch_table;
+
 template <typename T, typename RealT>
 static sycl::event heevd_batch_impl(sycl::queue &exec_q,
                                     const oneapi::mkl::job jobz,
@@ -175,8 +180,7 @@ void init_heevd_batch(py::module_ m)
                                   [dpctl_td_ns::num_types];
 
     {
-        evd::init_evd_dispatch_table<evd_batch_impl_fn_ptr_t,
-                                     HeevdBatchContigFactory>(
+        init_dispatch_table<evd_batch_impl_fn_ptr_t, HeevdBatchContigFactory>(
             heevd_batch_dispatch_table);
 
         auto heevd_batch_pyapi =
diff --git a/dpnp/backend/extensions/lapack/orgqr.cpp b/dpnp/backend/extensions/lapack/orgqr.cpp
index 33d67ac5a86a..296a4506ec0d 100644
--- a/dpnp/backend/extensions/lapack/orgqr.cpp
+++ b/dpnp/backend/extensions/lapack/orgqr.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 "orgqr.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 (*orgqr_impl_fn_ptr_t)(sycl::queue &,
                                            const std::int64_t,
                                            const std::int64_t,
@@ -251,9 +254,7 @@ struct OrgqrContigFactory
 
 void init_orgqr_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<orgqr_impl_fn_ptr_t, OrgqrContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(orgqr_dispatch_vector);
+    init_dispatch_vector<orgqr_impl_fn_ptr_t, OrgqrContigFactory>(
+        orgqr_dispatch_vector);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/lapack/orgqr_batch.cpp b/dpnp/backend/extensions/lapack/orgqr_batch.cpp
index c4689eabc34b..89324e2ba33f 100644
--- a/dpnp/backend/extensions/lapack/orgqr_batch.cpp
+++ b/dpnp/backend/extensions/lapack/orgqr_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/sycl_alloc_utils.hpp"
@@ -35,14 +38,14 @@
 #include "orgqr.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 (*orgqr_batch_impl_fn_ptr_t)(
     sycl::queue &,
     std::int64_t,
@@ -265,10 +268,7 @@ struct OrgqrBatchContigFactory
 
 void init_orgqr_batch_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<orgqr_batch_impl_fn_ptr_t,
-                                       OrgqrBatchContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(orgqr_batch_dispatch_vector);
+    init_dispatch_vector<orgqr_batch_impl_fn_ptr_t, OrgqrBatchContigFactory>(
+        orgqr_batch_dispatch_vector);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/lapack/potrf.cpp b/dpnp/backend/extensions/lapack/potrf.cpp
index fc6616f01cca..382a7dd8a770 100644
--- a/dpnp/backend/extensions/lapack/potrf.cpp
+++ b/dpnp/backend/extensions/lapack/potrf.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"
@@ -36,14 +39,14 @@
 #include "potrf.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 (*potrf_impl_fn_ptr_t)(sycl::queue &,
                                            const oneapi::mkl::uplo,
                                            const std::int64_t,
@@ -209,9 +212,7 @@ struct PotrfContigFactory
 
 void init_potrf_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<potrf_impl_fn_ptr_t, PotrfContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(potrf_dispatch_vector);
+    init_dispatch_vector<potrf_impl_fn_ptr_t, PotrfContigFactory>(
+        potrf_dispatch_vector);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/lapack/potrf_batch.cpp b/dpnp/backend/extensions/lapack/potrf_batch.cpp
index 1a56e4b13bde..6da47a960cc2 100644
--- a/dpnp/backend/extensions/lapack/potrf_batch.cpp
+++ b/dpnp/backend/extensions/lapack/potrf_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/sycl_alloc_utils.hpp"
@@ -36,14 +39,14 @@
 #include "potrf.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 (*potrf_batch_impl_fn_ptr_t)(
     sycl::queue &,
     const oneapi::mkl::uplo,
@@ -244,10 +247,7 @@ struct PotrfBatchContigFactory
 
 void init_potrf_batch_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<potrf_batch_impl_fn_ptr_t,
-                                       PotrfBatchContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(potrf_batch_dispatch_vector);
+    init_dispatch_vector<potrf_batch_impl_fn_ptr_t, PotrfBatchContigFactory>(
+        potrf_batch_dispatch_vector);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/lapack/syevd.cpp b/dpnp/backend/extensions/lapack/syevd.cpp
index 516589bf557a..98753119e680 100644
--- a/dpnp/backend/extensions/lapack/syevd.cpp
+++ b/dpnp/backend/extensions/lapack/syevd.cpp
@@ -30,6 +30,9 @@
 #include "evd_common.hpp"
 #include "syevd.hpp"
 
+// utils extension header
+#include "ext/common.hpp"
+
 // dpctl tensor headers
 #include "utils/type_utils.hpp"
 
@@ -38,6 +41,8 @@ namespace dpnp::extensions::lapack
 namespace mkl_lapack = oneapi::mkl::lapack;
 namespace type_utils = dpctl::tensor::type_utils;
 
+using ext::common::init_dispatch_table;
+
 template <typename T, typename RealT>
 static sycl::event syevd_impl(sycl::queue &exec_q,
                               const oneapi::mkl::job jobz,
@@ -137,7 +142,7 @@ void init_syevd(py::module_ m)
                                                  [dpctl_td_ns::num_types];
 
     {
-        evd::init_evd_dispatch_table<evd_impl_fn_ptr_t, SyevdContigFactory>(
+        init_dispatch_table<evd_impl_fn_ptr_t, SyevdContigFactory>(
             syevd_dispatch_table);
 
         auto syevd_pyapi = [&](sycl::queue &exec_q, const std::int8_t jobz,
diff --git a/dpnp/backend/extensions/lapack/syevd_batch.cpp b/dpnp/backend/extensions/lapack/syevd_batch.cpp
index 99ecd2911987..bff66b6527b3 100644
--- a/dpnp/backend/extensions/lapack/syevd_batch.cpp
+++ b/dpnp/backend/extensions/lapack/syevd_batch.cpp
@@ -31,6 +31,9 @@
 #include "evd_batch_common.hpp"
 #include "syevd_batch.hpp"
 
+// utils extension header
+#include "ext/common.hpp"
+
 // dpctl tensor headers
 #include "utils/type_utils.hpp"
 
@@ -39,6 +42,8 @@ namespace dpnp::extensions::lapack
 namespace mkl_lapack = oneapi::mkl::lapack;
 namespace type_utils = dpctl::tensor::type_utils;
 
+using ext::common::init_dispatch_table;
+
 template <typename T, typename RealT>
 static sycl::event syevd_batch_impl(sycl::queue &exec_q,
                                     const oneapi::mkl::job jobz,
@@ -175,8 +180,7 @@ void init_syevd_batch(py::module_ m)
                                   [dpctl_td_ns::num_types];
 
     {
-        evd::init_evd_dispatch_table<evd_batch_impl_fn_ptr_t,
-                                     SyevdBatchContigFactory>(
+        init_dispatch_table<evd_batch_impl_fn_ptr_t, SyevdBatchContigFactory>(
             syevd_batch_dispatch_table);
 
         auto syevd_batch_pyapi =
diff --git a/dpnp/backend/extensions/lapack/ungqr.cpp b/dpnp/backend/extensions/lapack/ungqr.cpp
index 73faae36d633..e23362310da5 100644
--- a/dpnp/backend/extensions/lapack/ungqr.cpp
+++ b/dpnp/backend/extensions/lapack/ungqr.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 "types_matrix.hpp"
 #include "ungqr.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 (*ungqr_impl_fn_ptr_t)(sycl::queue &,
                                            const std::int64_t,
                                            const std::int64_t,
@@ -251,9 +254,7 @@ struct UngqrContigFactory
 
 void init_ungqr_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<ungqr_impl_fn_ptr_t, UngqrContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(ungqr_dispatch_vector);
+    init_dispatch_vector<ungqr_impl_fn_ptr_t, UngqrContigFactory>(
+        ungqr_dispatch_vector);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/lapack/ungqr_batch.cpp b/dpnp/backend/extensions/lapack/ungqr_batch.cpp
index b5029bf2934d..a3ddc9bd5f1b 100644
--- a/dpnp/backend/extensions/lapack/ungqr_batch.cpp
+++ b/dpnp/backend/extensions/lapack/ungqr_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/sycl_alloc_utils.hpp"
@@ -35,14 +38,14 @@
 #include "types_matrix.hpp"
 #include "ungqr.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 (*ungqr_batch_impl_fn_ptr_t)(
     sycl::queue &,
     std::int64_t,
@@ -265,10 +268,7 @@ struct UngqrBatchContigFactory
 
 void init_ungqr_batch_dispatch_vector(void)
 {
-    dpctl_td_ns::DispatchVectorBuilder<ungqr_batch_impl_fn_ptr_t,
-                                       UngqrBatchContigFactory,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(ungqr_batch_dispatch_vector);
+    init_dispatch_vector<ungqr_batch_impl_fn_ptr_t, UngqrBatchContigFactory>(
+        ungqr_batch_dispatch_vector);
 }
 } // namespace dpnp::extensions::lapack
diff --git a/dpnp/backend/extensions/ufunc/CMakeLists.txt b/dpnp/backend/extensions/ufunc/CMakeLists.txt
index 1e8971b87c6f..5ee02ffe5fe1 100644
--- a/dpnp/backend/extensions/ufunc/CMakeLists.txt
+++ b/dpnp/backend/extensions/ufunc/CMakeLists.txt
@@ -27,7 +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/erf_funcs.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
diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp
index 6946b0af7f5a..287d36d13915 100644
--- a/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp
+++ b/dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp
@@ -27,7 +27,7 @@
 
 #include "bitwise_count.hpp"
 #include "degrees.hpp"
-#include "erf.hpp"
+#include "erf_funcs.hpp"
 #include "fabs.hpp"
 #include "fix.hpp"
 #include "float_power.hpp"
@@ -58,7 +58,7 @@ void init_elementwise_functions(py::module_ m)
 {
     init_bitwise_count(m);
     init_degrees(m);
-    init_erf(m);
+    init_erf_funcs(m);
     init_fabs(m);
     init_fix(m);
     init_float_power(m);
diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/erf.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/erf.cpp
deleted file mode 100644
index a0f3165a1083..000000000000
--- a/dpnp/backend/extensions/ufunc/elementwise_functions/erf.cpp
+++ /dev/null
@@ -1,132 +0,0 @@
-//*****************************************************************************
-// 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
diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/erf_funcs.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/erf_funcs.cpp
new file mode 100644
index 000000000000..1a441d23e830
--- /dev/null
+++ b/dpnp/backend/extensions/ufunc/elementwise_functions/erf_funcs.cpp
@@ -0,0 +1,232 @@
+//*****************************************************************************
+// 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_funcs.hpp"
+#include "kernels/elementwise_functions/erf.hpp"
+
+// utils extension header
+#include "ext/common.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;
+
+using ext::common::init_dispatch_vector;
+
+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
+ * an erf-like function in SYCL namespace 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;
+};
+
+static int output_typeid_vector[td_ns::num_types];
+
+template <typename fnT, typename T>
+struct TypeMapFactory
+{
+    std::enable_if_t<std::is_same<fnT, int>::value, int> get()
+    {
+        using rT = typename OutputType<T>::value_type;
+        return td_ns::GetTypeid<rT>{}.get();
+    }
+};
+
+using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
+using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
+
+#define MACRO_DEFINE_IMPL(__name__, __f_name__)                                \
+                                                                               \
+    using dpnp::kernels::erfs::__f_name__##Functor;                            \
+                                                                               \
+    template <typename argT, typename resT = argT, unsigned int vec_sz = 4,    \
+              unsigned int n_vecs = 2, bool enable_sg_loadstore = true>        \
+    using __f_name__##ContigFunctor =                                          \
+        ew_cmn_ns::UnaryContigFunctor<argT, resT,                              \
+                                      __f_name__##Functor<argT, resT>, vec_sz, \
+                                      n_vecs, enable_sg_loadstore>;            \
+                                                                               \
+    template <typename argTy, typename resTy, typename IndexerT>               \
+    using __f_name__##StridedFunctor =                                         \
+        ew_cmn_ns::UnaryStridedFunctor<argTy, resTy, IndexerT,                 \
+                                       __f_name__##Functor<argTy, resTy>>;     \
+                                                                               \
+    static unary_contig_impl_fn_ptr_t                                          \
+        __name__##_contig_dispatch_vector[td_ns::num_types];                   \
+    static unary_strided_impl_fn_ptr_t                                         \
+        __name__##_strided_dispatch_vector[td_ns::num_types];                  \
+                                                                               \
+    template <typename T1, typename T2, unsigned int vec_sz,                   \
+              unsigned int n_vecs>                                             \
+    class __name__##_contig_kernel;                                            \
+                                                                               \
+    template <typename argTy>                                                  \
+    sycl::event __name__##_contig_impl(                                        \
+        sycl::queue &exec_q, size_t nelems, const char *arg_p, char *res_p,    \
+        const std::vector<sycl::event> &depends = {})                          \
+    {                                                                          \
+        return ew_cmn_ns::unary_contig_impl<argTy, OutputType,                 \
+                                            __f_name__##ContigFunctor,         \
+                                            __name__##_contig_kernel>(         \
+            exec_q, nelems, arg_p, res_p, depends);                            \
+    }                                                                          \
+                                                                               \
+    template <typename fnT, typename T>                                        \
+    struct __f_name__##ContigFactory                                           \
+    {                                                                          \
+        fnT get()                                                              \
+        {                                                                      \
+            if constexpr (std::is_same_v<typename OutputType<T>::value_type,   \
+                                         void>) {                              \
+                fnT fn = nullptr;                                              \
+                return fn;                                                     \
+            }                                                                  \
+            else {                                                             \
+                fnT fn = __name__##_contig_impl<T>;                            \
+                return fn;                                                     \
+            }                                                                  \
+        }                                                                      \
+    };                                                                         \
+                                                                               \
+    template <typename T1, typename T2, typename T3>                           \
+    class __name__##_strided_kernel;                                           \
+                                                                               \
+    template <typename argTy>                                                  \
+    sycl::event __name__##_strided_impl(                                       \
+        sycl::queue &exec_q, size_t nelems, int nd,                            \
+        const py::ssize_t *shape_and_strides, const char *arg_p,               \
+        py::ssize_t arg_offset, char *res_p, py::ssize_t res_offset,           \
+        const std::vector<sycl::event> &depends,                               \
+        const std::vector<sycl::event> &additional_depends)                    \
+    {                                                                          \
+        return ew_cmn_ns::unary_strided_impl<argTy, OutputType,                \
+                                             __f_name__##StridedFunctor,       \
+                                             __name__##_strided_kernel>(       \
+            exec_q, nelems, nd, shape_and_strides, arg_p, arg_offset, res_p,   \
+            res_offset, depends, additional_depends);                          \
+    }                                                                          \
+                                                                               \
+    template <typename fnT, typename T>                                        \
+    struct __f_name__##StridedFactory                                          \
+    {                                                                          \
+        fnT get()                                                              \
+        {                                                                      \
+            if constexpr (std::is_same_v<typename OutputType<T>::value_type,   \
+                                         void>) {                              \
+                fnT fn = nullptr;                                              \
+                return fn;                                                     \
+            }                                                                  \
+            else {                                                             \
+                fnT fn = __name__##_strided_impl<T>;                           \
+                return fn;                                                     \
+            }                                                                  \
+        }                                                                      \
+    };
+
+template <template <typename fnT, typename T> typename contigFactoryT,
+          template <typename fnT, typename T>
+          typename stridedFactoryT>
+static void populate(py::module_ m,
+                     const char *name,
+                     const char *docstring,
+                     unary_contig_impl_fn_ptr_t *contig_dispatch_vector,
+                     unary_strided_impl_fn_ptr_t *strided_dispatch_vector)
+{
+    init_dispatch_vector<unary_contig_impl_fn_ptr_t, contigFactoryT>(
+        contig_dispatch_vector);
+
+    init_dispatch_vector<unary_strided_impl_fn_ptr_t, stridedFactoryT>(
+        strided_dispatch_vector);
+
+    using arrayT = dpctl::tensor::usm_ndarray;
+    auto pyapi = [&, contig_dispatch_vector, strided_dispatch_vector](
+                     const arrayT &src, const arrayT &dst, sycl::queue &exec_q,
+                     const std::vector<sycl::event> &depends = {}) {
+        return py_int::py_unary_ufunc(
+            src, dst, exec_q, depends, output_typeid_vector,
+            contig_dispatch_vector, strided_dispatch_vector);
+    };
+    m.def(name, pyapi, docstring, py::arg("sycl_queue"), py::arg("src"),
+          py::arg("dst"), py::arg("depends") = py::list());
+}
+
+MACRO_DEFINE_IMPL(erf, Erf);
+MACRO_DEFINE_IMPL(erfc, Erfc);
+} // namespace impl
+
+void init_erf_funcs(py::module_ m)
+{
+    using impl::output_typeid_vector;
+    init_dispatch_vector<int, impl::TypeMapFactory>(output_typeid_vector);
+
+    auto erf_result_type_pyapi = [&](const py::dtype &dtype) {
+        return py_int::py_unary_ufunc_result_type(dtype, output_typeid_vector);
+    };
+    m.def("_erf_result_type", erf_result_type_pyapi);
+
+    impl::populate<impl::ErfContigFactory, impl::ErfStridedFactory>(
+        m, "_erf", "", impl::erf_contig_dispatch_vector,
+        impl::erf_strided_dispatch_vector);
+
+    impl::populate<impl::ErfcContigFactory, impl::ErfcStridedFactory>(
+        m, "_erfc", "", impl::erfc_contig_dispatch_vector,
+        impl::erfc_strided_dispatch_vector);
+}
+} // namespace dpnp::extensions::ufunc
diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/erf.hpp b/dpnp/backend/extensions/ufunc/elementwise_functions/erf_funcs.hpp
similarity index 97%
rename from dpnp/backend/extensions/ufunc/elementwise_functions/erf.hpp
rename to dpnp/backend/extensions/ufunc/elementwise_functions/erf_funcs.hpp
index f02f7f4a6448..195f9d62da4f 100644
--- a/dpnp/backend/extensions/ufunc/elementwise_functions/erf.hpp
+++ b/dpnp/backend/extensions/ufunc/elementwise_functions/erf_funcs.hpp
@@ -31,5 +31,5 @@ namespace py = pybind11;
 
 namespace dpnp::extensions::ufunc
 {
-void init_erf(py::module_ m);
+void init_erf_funcs(py::module_ m);
 } // namespace dpnp::extensions::ufunc
diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/float_power.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/float_power.cpp
index 3ce168940148..6abf2f217a11 100644
--- a/dpnp/backend/extensions/ufunc/elementwise_functions/float_power.cpp
+++ b/dpnp/backend/extensions/ufunc/elementwise_functions/float_power.cpp
@@ -33,6 +33,9 @@
 
 #include "float_power.hpp"
 
+// utils extension header
+#include "ext/common.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
@@ -50,6 +53,8 @@ namespace impl
 {
 namespace td_ns = dpctl::tensor::type_dispatch;
 
+using ext::common::init_dispatch_table;
+
 // Supports only float and complex types
 template <typename T1, typename T2>
 struct OutputType
@@ -82,10 +87,9 @@ struct TypeMapFactory
     }
 };
 
-void populate_float_power_dispatch_tables(void)
+static void populate_float_power_dispatch_tables(void)
 {
-    td_ns::DispatchTableBuilder<int, TypeMapFactory, td_ns::num_types> dvb;
-    dvb.populate_dispatch_table(float_power_output_typeid_table);
+    init_dispatch_table<int, TypeMapFactory>(float_power_output_typeid_table);
 }
 } // namespace impl
 
diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/interpolate.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/interpolate.cpp
index d88b255e6271..ad0955f07f35 100644
--- a/dpnp/backend/extensions/ufunc/elementwise_functions/interpolate.cpp
+++ b/dpnp/backend/extensions/ufunc/elementwise_functions/interpolate.cpp
@@ -44,6 +44,7 @@
 
 #include "kernels/elementwise_functions/interpolate.hpp"
 
+// utils extension headers
 #include "ext/common.hpp"
 #include "ext/validation_utils.hpp"
 
@@ -67,6 +68,7 @@ namespace dpnp::extensions::ufunc
 
 namespace impl
 {
+using ext::common::init_dispatch_vector;
 
 template <typename T>
 using value_type_of_t = typename value_type_of<T>::type;
@@ -242,13 +244,10 @@ struct InterpolateFactory
     }
 };
 
-void init_interpolate_dispatch_vectors()
+static void init_interpolate_dispatch_vectors()
 {
-    using namespace td_ns;
-
-    DispatchVectorBuilder<interpolate_fn_ptr_t, InterpolateFactory, num_types>
-        dtb_interpolate;
-    dtb_interpolate.populate_dispatch_vector(interpolate_dispatch_vector);
+    init_dispatch_vector<interpolate_fn_ptr_t, InterpolateFactory>(
+        interpolate_dispatch_vector);
 }
 
 } // namespace impl
diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/isclose.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/isclose.cpp
index 43ff66a9a9cb..b96ecd365909 100644
--- a/dpnp/backend/extensions/ufunc/elementwise_functions/isclose.cpp
+++ b/dpnp/backend/extensions/ufunc/elementwise_functions/isclose.cpp
@@ -45,6 +45,7 @@
 #include "utils/type_dispatch.hpp"
 #include "utils/type_utils.hpp"
 
+// utils extension headers
 #include "ext/common.hpp"
 #include "ext/validation_utils.hpp"
 
@@ -68,6 +69,7 @@ namespace dpnp::extensions::ufunc
 
 namespace impl
 {
+using ext::common::init_dispatch_vector;
 
 typedef sycl::event (*isclose_strided_scalar_fn_ptr_t)(
     sycl::queue &,
@@ -351,19 +353,14 @@ struct IsCloseContigScalarFactory
     }
 };
 
-void populate_isclose_dispatch_vectors()
+static void populate_isclose_dispatch_vectors()
 {
-    using namespace td_ns;
-
-    DispatchVectorBuilder<isclose_strided_scalar_fn_ptr_t,
-                          IsCloseStridedScalarFactory, num_types>
-        dvb1;
-    dvb1.populate_dispatch_vector(isclose_strided_scalar_dispatch_vector);
-
-    DispatchVectorBuilder<isclose_contig_scalar_fn_ptr_t,
-                          IsCloseContigScalarFactory, num_types>
-        dvb2;
-    dvb2.populate_dispatch_vector(isclose_contig_dispatch_vector);
+    init_dispatch_vector<isclose_strided_scalar_fn_ptr_t,
+                         IsCloseStridedScalarFactory>(
+        isclose_strided_scalar_dispatch_vector);
+    init_dispatch_vector<isclose_contig_scalar_fn_ptr_t,
+                         IsCloseContigScalarFactory>(
+        isclose_contig_dispatch_vector);
 }
 
 } // namespace impl
diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/nan_to_num.cpp b/dpnp/backend/extensions/ufunc/elementwise_functions/nan_to_num.cpp
index fae8f8bc331c..0a700bbd695e 100644
--- a/dpnp/backend/extensions/ufunc/elementwise_functions/nan_to_num.cpp
+++ b/dpnp/backend/extensions/ufunc/elementwise_functions/nan_to_num.cpp
@@ -52,6 +52,7 @@
 #include "utils/type_dispatch.hpp"
 #include "utils/type_utils.hpp"
 
+// utils extension header
 #include "ext/common.hpp"
 
 namespace py = pybind11;
@@ -65,6 +66,7 @@ namespace dpnp::extensions::ufunc
 
 namespace impl
 {
+using ext::common::init_dispatch_vector;
 
 template <typename T>
 using value_type_of_t = typename value_type_of<T>::type;
@@ -370,17 +372,12 @@ struct NanToNumContigFactory
     }
 };
 
-void populate_nan_to_num_dispatch_vectors(void)
+static void populate_nan_to_num_dispatch_vectors(void)
 {
-    using namespace td_ns;
-
-    DispatchVectorBuilder<nan_to_num_fn_ptr_t, NanToNumFactory, num_types> dvb1;
-    dvb1.populate_dispatch_vector(nan_to_num_dispatch_vector);
-
-    DispatchVectorBuilder<nan_to_num_contig_fn_ptr_t, NanToNumContigFactory,
-                          num_types>
-        dvb2;
-    dvb2.populate_dispatch_vector(nan_to_num_contig_dispatch_vector);
+    init_dispatch_vector<nan_to_num_fn_ptr_t, NanToNumFactory>(
+        nan_to_num_dispatch_vector);
+    init_dispatch_vector<nan_to_num_contig_fn_ptr_t, NanToNumContigFactory>(
+        nan_to_num_contig_dispatch_vector);
 }
 
 } // namespace impl
diff --git a/dpnp/backend/extensions/ufunc/elementwise_functions/populate.hpp b/dpnp/backend/extensions/ufunc/elementwise_functions/populate.hpp
index d0ffe8abfd9f..98a4b2702cc2 100644
--- a/dpnp/backend/extensions/ufunc/elementwise_functions/populate.hpp
+++ b/dpnp/backend/extensions/ufunc/elementwise_functions/populate.hpp
@@ -25,6 +25,11 @@
 
 #pragma once
 
+// utils extension header
+#include "ext/common.hpp"
+
+namespace ext_ns = ext::common;
+
 /**
  * @brief A macro used to define factories and a populating unary universal
  * functions.
@@ -107,19 +112,14 @@
                                                                                \
     void populate_##__name__##_dispatch_vectors(void)                          \
     {                                                                          \
-        td_ns::DispatchVectorBuilder<unary_contig_impl_fn_ptr_t,               \
-                                     ContigFactory, td_ns::num_types>          \
-            dvb1;                                                              \
-        dvb1.populate_dispatch_vector(__name__##_contig_dispatch_vector);      \
-                                                                               \
-        td_ns::DispatchVectorBuilder<unary_strided_impl_fn_ptr_t,              \
-                                     StridedFactory, td_ns::num_types>         \
-            dvb2;                                                              \
-        dvb2.populate_dispatch_vector(__name__##_strided_dispatch_vector);     \
-                                                                               \
-        td_ns::DispatchVectorBuilder<int, TypeMapFactory, td_ns::num_types>    \
-            dvb3;                                                              \
-        dvb3.populate_dispatch_vector(__name__##_output_typeid_vector);        \
+        ext_ns::init_dispatch_vector<unary_contig_impl_fn_ptr_t,               \
+                                     ContigFactory>(                           \
+            __name__##_contig_dispatch_vector);                                \
+        ext_ns::init_dispatch_vector<unary_strided_impl_fn_ptr_t,              \
+                                     StridedFactory>(                          \
+            __name__##_strided_dispatch_vector);                               \
+        ext_ns::init_dispatch_vector<int, TypeMapFactory>(                     \
+            __name__##_output_typeid_vector);                                  \
     };
 
 /**
@@ -214,17 +214,12 @@
                                                                                \
     void populate_##__name__##_dispatch_tables(void)                           \
     {                                                                          \
-        td_ns::DispatchTableBuilder<binary_contig_impl_fn_ptr_t,               \
-                                    ContigFactory, td_ns::num_types>           \
-            dvb1;                                                              \
-        dvb1.populate_dispatch_table(__name__##_contig_dispatch_table);        \
-                                                                               \
-        td_ns::DispatchTableBuilder<binary_strided_impl_fn_ptr_t,              \
-                                    StridedFactory, td_ns::num_types>          \
-            dvb2;                                                              \
-        dvb2.populate_dispatch_table(__name__##_strided_dispatch_table);       \
-                                                                               \
-        td_ns::DispatchTableBuilder<int, TypeMapFactory, td_ns::num_types>     \
-            dvb3;                                                              \
-        dvb3.populate_dispatch_table(__name__##_output_typeid_table);          \
+        ext_ns::init_dispatch_table<binary_contig_impl_fn_ptr_t,               \
+                                    ContigFactory>(                            \
+            __name__##_contig_dispatch_table);                                 \
+        ext_ns::init_dispatch_table<binary_strided_impl_fn_ptr_t,              \
+                                    StridedFactory>(                           \
+            __name__##_strided_dispatch_table);                                \
+        ext_ns::init_dispatch_table<int, TypeMapFactory>(                      \
+            __name__##_output_typeid_table);                                   \
     };
diff --git a/dpnp/backend/extensions/vm/CMakeLists.txt b/dpnp/backend/extensions/vm/CMakeLists.txt
index 057311c146c9..1c12f404bca8 100644
--- a/dpnp/backend/extensions/vm/CMakeLists.txt
+++ b/dpnp/backend/extensions/vm/CMakeLists.txt
@@ -42,7 +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}/erf_funcs.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/exp.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/exp2.cpp
         ${CMAKE_CURRENT_SOURCE_DIR}/expm1.cpp
@@ -95,8 +95,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})
diff --git a/dpnp/backend/extensions/vm/common.hpp b/dpnp/backend/extensions/vm/common.hpp
index 3a93d7b6e8e9..16a166a68bff 100644
--- a/dpnp/backend/extensions/vm/common.hpp
+++ b/dpnp/backend/extensions/vm/common.hpp
@@ -33,15 +33,27 @@
 #include <dpctl4pybind11.hpp>
 #include <pybind11/pybind11.h>
 
+// utils extension header
+#include "ext/common.hpp"
+
 // dpctl tensor headers
 #include "utils/memory_overlap.hpp"
 #include "utils/type_dispatch.hpp"
 
-#include "dpnp_utils.hpp"
+/**
+ * Version of Intel MKL at which transition to OneMKL release 2023.2.0 occurs.
+ *
+ * @note with OneMKL=2023.1.0 the call of oneapi::mkl::vm::div() was dead
+ * locked inside ~usm_wrapper_to_host()->{...; q_->wait_and_throw(); ...}
+ */
+#ifndef __INTEL_MKL_2023_2_0_VERSION_REQUIRED
+#define __INTEL_MKL_2023_2_0_VERSION_REQUIRED 20230002L
+#endif
 
 static_assert(INTEL_MKL_VERSION >= __INTEL_MKL_2023_2_0_VERSION_REQUIRED,
               "OneMKL does not meet minimum version requirement");
 
+namespace ext_ns = ext::common;
 namespace py = pybind11;
 namespace td_ns = dpctl::tensor::type_dispatch;
 
@@ -278,11 +290,10 @@ bool need_to_call_binary_ufunc(sycl::queue &exec_q,
                                                                                \
     static void populate_dispatch_vectors(void)                                \
     {                                                                          \
-        py_internal::init_ufunc_dispatch_vector<int, TypeMapFactory>(          \
+        ext_ns::init_dispatch_vector<int, TypeMapFactory>(                     \
             output_typeid_vector);                                             \
-        py_internal::init_ufunc_dispatch_vector<unary_contig_impl_fn_ptr_t,    \
-                                                ContigFactory>(                \
-            contig_dispatch_vector);                                           \
+        ext_ns::init_dispatch_vector<unary_contig_impl_fn_ptr_t,               \
+                                     ContigFactory>(contig_dispatch_vector);   \
     };
 
 /**
@@ -319,30 +330,9 @@ bool need_to_call_binary_ufunc(sycl::queue &exec_q,
                                                                                \
     static void populate_dispatch_tables(void)                                 \
     {                                                                          \
-        py_internal::init_ufunc_dispatch_table<int, TypeMapFactory>(           \
+        ext_ns::init_dispatch_table<int, TypeMapFactory>(                      \
             output_typeid_vector);                                             \
-        py_internal::init_ufunc_dispatch_table<binary_contig_impl_fn_ptr_t,    \
-                                               ContigFactory>(                 \
-            contig_dispatch_vector);                                           \
+        ext_ns::init_dispatch_table<binary_contig_impl_fn_ptr_t,               \
+                                    ContigFactory>(contig_dispatch_vector);    \
     };
-
-template <typename dispatchT,
-          template <typename fnT, typename T>
-          typename factoryT,
-          int _num_types = td_ns::num_types>
-void init_ufunc_dispatch_vector(dispatchT dispatch_vector[])
-{
-    td_ns::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 = td_ns::num_types>
-void init_ufunc_dispatch_table(dispatchT dispatch_table[][_num_types])
-{
-    td_ns::DispatchTableBuilder<dispatchT, factoryT, _num_types> dtb;
-    dtb.populate_dispatch_table(dispatch_table);
-}
 } // namespace dpnp::extensions::vm::py_internal
diff --git a/dpnp/backend/extensions/vm/erf.cpp b/dpnp/backend/extensions/vm/erf.cpp
deleted file mode 100644
index 23a396b8bc64..000000000000
--- a/dpnp/backend/extensions/vm/erf.cpp
+++ /dev/null
@@ -1,140 +0,0 @@
-//*****************************************************************************
-// 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 <cstddef>
-#include <cstdint>
-#include <type_traits>
-#include <vector>
-
-#include <oneapi/mkl.hpp>
-#include <sycl/sycl.hpp>
-
-#include "dpctl4pybind11.hpp"
-
-#include "common.hpp"
-#include "erf.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"
-#include "utils/type_utils.hpp"
-
-namespace dpnp::extensions::vm
-{
-namespace py = pybind11;
-namespace py_int = dpnp::extensions::py_internal;
-namespace td_ns = dpctl::tensor::type_dispatch;
-
-namespace impl
-{
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
-namespace mkl_vm = oneapi::mkl::vm; // OneMKL namespace with VM functions
-namespace tu_ns = dpctl::tensor::type_utils;
-
-/**
- * @brief A factory to define pairs of supported types for which
- * MKL VM library provides support in oneapi::mkl::vm::erf<T> function.
- *
- * @tparam T Type of input vector `a` and of result vector `y`.
- */
-template <typename T>
-struct OutputType
-{
-    using value_type =
-        typename std::disjunction<td_ns::TypeMapResultEntry<T, float>,
-                                  td_ns::TypeMapResultEntry<T, double>,
-                                  td_ns::DefaultResultEntry<void>>::result_type;
-};
-
-template <typename T>
-static sycl::event erf_contig_impl(sycl::queue &exec_q,
-                                   std::size_t in_n,
-                                   const char *in_a,
-                                   char *out_y,
-                                   const std::vector<sycl::event> &depends)
-{
-    tu_ns::validate_type_for_device<T>(exec_q);
-
-    std::int64_t n = static_cast<std::int64_t>(in_n);
-    const T *a = reinterpret_cast<const T *>(in_a);
-
-    using resTy = typename OutputType<T>::value_type;
-    resTy *y = reinterpret_cast<resTy *>(out_y);
-
-    return mkl_vm::erf(exec_q,
-                       n, // number of elements to be calculated
-                       a, // pointer `a` containing input vector of size n
-                       y, // pointer `y` to the output vector of size n
-                       depends);
-}
-
-using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
-using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
-
-static int output_typeid_vector[td_ns::num_types];
-static unary_contig_impl_fn_ptr_t contig_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_dispatch_vectors();
-    using impl::contig_dispatch_vector;
-    using impl::output_typeid_vector;
-
-    auto erf_pyapi = [&](sycl::queue &exec_q, const arrayT &src,
-                         const arrayT &dst, const event_vecT &depends = {}) {
-        return py_int::py_unary_ufunc(
-            src, dst, exec_q, depends, output_typeid_vector,
-            contig_dispatch_vector,
-            // no support of strided implementation in OneMKL
-            td_ns::NullPtrVector<impl::unary_strided_impl_fn_ptr_t>{});
-    };
-    m.def("_erf", erf_pyapi,
-          "Call `erf` function from OneMKL VM library to compute "
-          "the natural (base-e) erfonential of vector elements",
-          py::arg("sycl_queue"), py::arg("src"), py::arg("dst"),
-          py::arg("depends") = py::list());
-
-    auto erf_need_to_call_pyapi = [&](sycl::queue &exec_q, const arrayT &src,
-                                      const arrayT &dst) {
-        return py_internal::need_to_call_unary_ufunc(
-            exec_q, src, dst, output_typeid_vector, contig_dispatch_vector);
-    };
-    m.def("_mkl_erf_to_call", erf_need_to_call_pyapi,
-          "Check input arguments to answer if `erf` function from "
-          "OneMKL VM library can be used",
-          py::arg("sycl_queue"), py::arg("src"), py::arg("dst"));
-}
-} // namespace dpnp::extensions::vm
diff --git a/dpnp/backend/extensions/vm/erf_funcs.cpp b/dpnp/backend/extensions/vm/erf_funcs.cpp
new file mode 100644
index 000000000000..2414c56a9e50
--- /dev/null
+++ b/dpnp/backend/extensions/vm/erf_funcs.cpp
@@ -0,0 +1,188 @@
+//*****************************************************************************
+// 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 <cstddef>
+#include <cstdint>
+#include <type_traits>
+#include <vector>
+
+#include <oneapi/mkl.hpp>
+#include <sycl/sycl.hpp>
+
+#include "dpctl4pybind11.hpp"
+
+#include "common.hpp"
+#include "erf_funcs.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"
+#include "utils/type_utils.hpp"
+
+namespace dpnp::extensions::vm
+{
+namespace py = pybind11;
+namespace py_int = dpnp::extensions::py_internal;
+namespace td_ns = dpctl::tensor::type_dispatch;
+
+using ext::common::init_dispatch_vector;
+
+namespace impl
+{
+namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
+namespace mkl_vm = oneapi::mkl::vm; // OneMKL namespace with VM functions
+namespace tu_ns = dpctl::tensor::type_utils;
+
+/**
+ * @brief A factory to define pairs of supported types for which
+ * MKL VM library provides support for a family of erf-like functions.
+ *
+ * @tparam T Type of input vector `a` and of result vector `y`.
+ */
+template <typename T>
+struct OutputType
+{
+    using value_type =
+        typename std::disjunction<td_ns::TypeMapResultEntry<T, float>,
+                                  td_ns::TypeMapResultEntry<T, double>,
+                                  td_ns::DefaultResultEntry<void>>::result_type;
+};
+
+static int output_typeid_vector[td_ns::num_types];
+
+template <typename fnT, typename T>
+struct TypeMapFactory
+{
+    std::enable_if_t<std::is_same<fnT, int>::value, int> get()
+    {
+        using rT = typename OutputType<T>::value_type;
+        return td_ns::GetTypeid<rT>{}.get();
+    }
+};
+
+using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
+
+#define MACRO_DEFINE_IMPL(__name__, __f_name__)                                \
+    template <typename T>                                                      \
+    static sycl::event __name__##_contig_impl(                                 \
+        sycl::queue &exec_q, std::size_t in_n, const char *in_a, char *out_y,  \
+        const std::vector<sycl::event> &depends)                               \
+    {                                                                          \
+        tu_ns::validate_type_for_device<T>(exec_q);                            \
+                                                                               \
+        std::int64_t n = static_cast<std::int64_t>(in_n);                      \
+        const T *a = reinterpret_cast<const T *>(in_a);                        \
+                                                                               \
+        using resTy = typename OutputType<T>::value_type;                      \
+        resTy *y = reinterpret_cast<resTy *>(out_y);                           \
+                                                                               \
+        return mkl_vm::__name__(                                               \
+            exec_q, n, /* number of elements to be calculated*/                \
+            a,         /* pointer `a` containing input vector of size n*/      \
+            y,         /* pointer `y` to the output vector of size n*/         \
+            depends);                                                          \
+    }                                                                          \
+                                                                               \
+    static unary_contig_impl_fn_ptr_t                                          \
+        __name__##_contig_dispatch_vector[td_ns::num_types];                   \
+                                                                               \
+    template <typename fnT, typename T>                                        \
+    struct __f_name__##ContigFactory                                           \
+    {                                                                          \
+        fnT get()                                                              \
+        {                                                                      \
+            if constexpr (std::is_same_v<typename OutputType<T>::value_type,   \
+                                         void>) {                              \
+                return nullptr;                                                \
+            }                                                                  \
+            else {                                                             \
+                return __name__##_contig_impl<T>;                              \
+            }                                                                  \
+        }                                                                      \
+    };
+
+MACRO_DEFINE_IMPL(erf, Erf);
+MACRO_DEFINE_IMPL(erfc, Erfc);
+
+template <template <typename fnT, typename T> typename factoryT>
+static void populate(py::module_ m,
+                     const char *name,
+                     const char *docstring,
+                     unary_contig_impl_fn_ptr_t *contig_dispatch_vector)
+{
+    init_dispatch_vector<unary_contig_impl_fn_ptr_t, factoryT>(
+        contig_dispatch_vector);
+
+    using arrayT = dpctl::tensor::usm_ndarray;
+    auto pyapi = [&, contig_dispatch_vector](
+                     sycl::queue &exec_q, const arrayT &src, const arrayT &dst,
+                     const std::vector<sycl::event> &depends = {}) {
+        return py_int::py_unary_ufunc(
+            src, dst, exec_q, depends, output_typeid_vector,
+            contig_dispatch_vector,
+            // no support of strided implementation in OneMKL
+            td_ns::NullPtrVector<ew_cmn_ns::unary_strided_impl_fn_ptr_t>{});
+    };
+    m.def(name, pyapi, docstring, py::arg("sycl_queue"), py::arg("src"),
+          py::arg("dst"), py::arg("depends") = py::list());
+}
+} // namespace impl
+
+void init_erf_funcs(py::module_ m)
+{
+    using arrayT = dpctl::tensor::usm_ndarray;
+    using impl::output_typeid_vector;
+
+    init_dispatch_vector<int, impl::TypeMapFactory>(output_typeid_vector);
+
+    auto erf_need_to_call_pyapi = [&](sycl::queue &exec_q, const arrayT &src,
+                                      const arrayT &dst) {
+        return py_internal::need_to_call_unary_ufunc(
+            exec_q, src, dst, output_typeid_vector,
+            impl::erf_contig_dispatch_vector);
+    };
+    m.def("_mkl_erf_to_call", erf_need_to_call_pyapi,
+          "Check input arguments to answer if any erf-like function from "
+          "OneMKL VM library can be used",
+          py::arg("sycl_queue"), py::arg("src"), py::arg("dst"));
+
+    impl::populate<impl::ErfContigFactory>(
+        m, "_erf",
+        "Call `erf` function from OneMKL VM library to compute the error "
+        "function value of vector elements",
+        impl::erf_contig_dispatch_vector);
+
+    impl::populate<impl::ErfcContigFactory>(
+        m, "_erfc",
+        "Call `erfc` function from OneMKL VM library to compute the "
+        "complementary error function value of vector elements",
+        impl::erfc_contig_dispatch_vector);
+}
+} // namespace dpnp::extensions::vm
diff --git a/dpnp/backend/extensions/vm/erf.hpp b/dpnp/backend/extensions/vm/erf_funcs.hpp
similarity index 97%
rename from dpnp/backend/extensions/vm/erf.hpp
rename to dpnp/backend/extensions/vm/erf_funcs.hpp
index 25eb2b05a7cc..771ba474aa49 100644
--- a/dpnp/backend/extensions/vm/erf.hpp
+++ b/dpnp/backend/extensions/vm/erf_funcs.hpp
@@ -31,5 +31,5 @@ namespace py = pybind11;
 
 namespace dpnp::extensions::vm
 {
-void init_erf(py::module_ m);
+void init_erf_funcs(py::module_ m);
 } // namespace dpnp::extensions::vm
diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp
index 33015179b73c..54b7cbab2318 100644
--- a/dpnp/backend/extensions/vm/vm_py.cpp
+++ b/dpnp/backend/extensions/vm/vm_py.cpp
@@ -45,7 +45,7 @@
 #include "cos.hpp"
 #include "cosh.hpp"
 #include "div.hpp"
-#include "erf.hpp"
+#include "erf_funcs.hpp"
 #include "exp.hpp"
 #include "exp2.hpp"
 #include "expm1.hpp"
@@ -98,7 +98,7 @@ PYBIND11_MODULE(_vm_impl, m)
     vm_ns::init_cos(m);
     vm_ns::init_cosh(m);
     vm_ns::init_div(m);
-    vm_ns::init_erf(m);
+    vm_ns::init_erf_funcs(m);
     vm_ns::init_exp(m);
     vm_ns::init_exp2(m);
     vm_ns::init_expm1(m);
diff --git a/dpnp/backend/extensions/window/CMakeLists.txt b/dpnp/backend/extensions/window/CMakeLists.txt
index 27b60c0fdc64..f8bd38d3a469 100644
--- a/dpnp/backend/extensions/window/CMakeLists.txt
+++ b/dpnp/backend/extensions/window/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})
diff --git a/dpnp/backend/extensions/window/common.hpp b/dpnp/backend/extensions/window/common.hpp
index 9a18d135b012..584e5aeaa639 100644
--- a/dpnp/backend/extensions/window/common.hpp
+++ b/dpnp/backend/extensions/window/common.hpp
@@ -30,6 +30,8 @@
 #include <sycl/sycl.hpp>
 
 #include "dpctl4pybind11.hpp"
+
+// dpctl tensor headers
 #include "utils/output_validation.hpp"
 #include "utils/type_dispatch.hpp"
 #include "utils/type_utils.hpp"
@@ -127,18 +129,4 @@ inline std::pair<sycl::event, sycl::event>
 
     return std::make_pair(args_ev, window_ev);
 }
-
-template <typename funcPtrT,
-          template <typename fnT, typename T>
-          typename factoryT>
-void init_window_dispatch_vectors(funcPtrT window_dispatch_vector[])
-{
-    dpctl_td_ns::DispatchVectorBuilder<funcPtrT, factoryT,
-                                       dpctl_td_ns::num_types>
-        contig;
-    contig.populate_dispatch_vector(window_dispatch_vector);
-
-    return;
-}
-
 } // namespace dpnp::extensions::window
diff --git a/dpnp/backend/extensions/window/kaiser.cpp b/dpnp/backend/extensions/window/kaiser.cpp
index 33f5551b32e3..ade5f6450fa7 100644
--- a/dpnp/backend/extensions/window/kaiser.cpp
+++ b/dpnp/backend/extensions/window/kaiser.cpp
@@ -26,18 +26,24 @@
 #include "kaiser.hpp"
 #include "common.hpp"
 
+// utils extension header
+#include "ext/common.hpp"
+
+// dpctl tensor headers
 #include "utils/output_validation.hpp"
 #include "utils/type_dispatch.hpp"
 #include "utils/type_utils.hpp"
 
 #include <sycl/sycl.hpp>
 
-#include "../kernels/elementwise_functions/i0.hpp"
+#include "kernels/elementwise_functions/i0.hpp"
 
 namespace dpnp::extensions::window
 {
 namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
 
+using ext::common::init_dispatch_vector;
+
 typedef sycl::event (*kaiser_fn_ptr_t)(sycl::queue &,
                                        char *,
                                        const std::size_t,
@@ -132,7 +138,7 @@ std::pair<sycl::event, sycl::event>
 
 void init_kaiser_dispatch_vectors()
 {
-    init_window_dispatch_vectors<kaiser_fn_ptr_t, KaiserFactory>(
+    init_dispatch_vector<kaiser_fn_ptr_t, KaiserFactory>(
         kaiser_dispatch_vector);
 }
 
diff --git a/dpnp/backend/extensions/window/window_py.cpp b/dpnp/backend/extensions/window/window_py.cpp
index 3324fe98e1d2..94830183b4f1 100644
--- a/dpnp/backend/extensions/window/window_py.cpp
+++ b/dpnp/backend/extensions/window/window_py.cpp
@@ -37,8 +37,13 @@
 #include "hanning.hpp"
 #include "kaiser.hpp"
 
+// utils extension header
+#include "ext/common.hpp"
+
 namespace window_ns = dpnp::extensions::window;
 namespace py = pybind11;
+
+using ext::common::init_dispatch_vector;
 using window_ns::window_fn_ptr_t;
 
 namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
@@ -54,8 +59,8 @@ PYBIND11_MODULE(_window_impl, m)
     using event_vecT = std::vector<sycl::event>;
 
     {
-        window_ns::init_window_dispatch_vectors<
-            window_ns::window_fn_ptr_t, window_ns::kernels::BartlettFactory>(
+        init_dispatch_vector<window_ns::window_fn_ptr_t,
+                             window_ns::kernels::BartlettFactory>(
             bartlett_dispatch_vector);
 
         auto bartlett_pyapi = [&](sycl::queue &exec_q, const arrayT &result,
@@ -70,8 +75,8 @@ PYBIND11_MODULE(_window_impl, m)
     }
 
     {
-        window_ns::init_window_dispatch_vectors<
-            window_ns::window_fn_ptr_t, window_ns::kernels::BlackmanFactory>(
+        init_dispatch_vector<window_ns::window_fn_ptr_t,
+                             window_ns::kernels::BlackmanFactory>(
             blackman_dispatch_vector);
 
         auto blackman_pyapi = [&](sycl::queue &exec_q, const arrayT &result,
@@ -86,8 +91,8 @@ PYBIND11_MODULE(_window_impl, m)
     }
 
     {
-        window_ns::init_window_dispatch_vectors<
-            window_ns::window_fn_ptr_t, window_ns::kernels::HammingFactory>(
+        init_dispatch_vector<window_ns::window_fn_ptr_t,
+                             window_ns::kernels::HammingFactory>(
             hamming_dispatch_vector);
 
         auto hamming_pyapi = [&](sycl::queue &exec_q, const arrayT &result,
@@ -102,8 +107,8 @@ PYBIND11_MODULE(_window_impl, m)
     }
 
     {
-        window_ns::init_window_dispatch_vectors<
-            window_ns::window_fn_ptr_t, window_ns::kernels::HanningFactory>(
+        init_dispatch_vector<window_ns::window_fn_ptr_t,
+                             window_ns::kernels::HanningFactory>(
             hanning_dispatch_vector);
 
         auto hanning_pyapi = [&](sycl::queue &exec_q, const arrayT &result,
diff --git a/dpnp/backend/include/dpnp_gen_1arg_1type_tbl.hpp b/dpnp/backend/include/dpnp_gen_1arg_1type_tbl.hpp
deleted file mode 100644
index 8a415a57bf8f..000000000000
--- a/dpnp/backend/include/dpnp_gen_1arg_1type_tbl.hpp
+++ /dev/null
@@ -1,94 +0,0 @@
-//*****************************************************************************
-// Copyright (c) 2016-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.
-//*****************************************************************************
-
-#if defined(MACRO_1ARG_1TYPE_OP)
-
-/*
- * This header file contains single argument element wise functions definitions
- *
- * Macro `MACRO_1ARG_1TYPE_OP` must be defined before usage
- *
- * Parameters:
- * - public name of the function and kernel name
- * - operation used to calculate the result
- * - mkl operation used to calculate the result
- *
- */
-
-#ifdef _SECTION_DOCUMENTATION_GENERATION_
-
-#define MACRO_1ARG_1TYPE_OP(__name__, __operation1__, __operation2__)          \
-    /** @ingroup BACKEND_API */                                                \
-    /** @brief Per element operation function __name__ */                      \
-    /** */                                                                     \
-    /** Function "__name__" executes operator "__operation1__" over            \
-     * each element of the array                        */                     \
-    /** */                                                                     \
-    /** @param[in]  q_ref              Reference to SYCL queue. */             \
-    /** @param[out] result_out         Output array. */                        \
-    /** @param[in]  result_size        Output array size. */                   \
-    /** @param[in]  result_ndim        Number of output array                  \
-     * dimensions. */                                                          \
-    /** @param[in]  result_shape       Output array shape. */                  \
-    /** @param[in]  result_strides     Output array strides. */                \
-    /** @param[in]  input1_in          Input array 1. */                       \
-    /** @param[in]  input1_size        Input array 1 size. */                  \
-    /** @param[in]  input1_ndim        Number of input array 1                 \
-     * dimensions. */                                                          \
-    /** @param[in]  input1_shape       Input array 1 shape. */                 \
-    /** @param[in]  input1_strides     Input array 1 strides. */               \
-    /** @param[in]  where              Where condition. */                     \
-    /** @param[in]  dep_event_vec_ref  Reference to vector of SYCL             \
-     * events. */                                                              \
-    template <typename _DataType>                                              \
-    DPCTLSyclEventRef __name__(                                                \
-        DPCTLSyclQueueRef q_ref, void *result_out, const size_t result_size,   \
-        const size_t result_ndim, const shape_elem_type *result_shape,         \
-        const shape_elem_type *result_strides, const void *input1_in,          \
-        const size_t input1_size, const size_t input1_ndim,                    \
-        const shape_elem_type *input1_shape,                                   \
-        const shape_elem_type *input1_strides, const size_t *where,            \
-        const DPCTLEventVectorRef dep_event_vec_ref);                          \
-                                                                               \
-    template <typename _DataType>                                              \
-    void __name__(                                                             \
-        void *result_out, const size_t result_size, const size_t result_ndim,  \
-        const shape_elem_type *result_shape,                                   \
-        const shape_elem_type *result_strides, const void *input1_in,          \
-        const size_t input1_size, const size_t input1_ndim,                    \
-        const shape_elem_type *input1_shape,                                   \
-        const shape_elem_type *input1_strides, const size_t *where);
-
-#endif // _SECTION_DOCUMENTATION_GENERATION_
-
-MACRO_1ARG_1TYPE_OP(dpnp_erf_c,
-                    dispatch_erf_op(input_elem),
-                    oneapi::mkl::vm::erf(q, input1_size, input1_data, result))
-
-#undef MACRO_1ARG_1TYPE_OP
-
-#else
-#error "MACRO_1ARG_1TYPE_OP is not defined"
-#endif // MACRO_1ARG_1TYPE_OP
diff --git a/dpnp/backend/kernels/elementwise_functions/erf.hpp b/dpnp/backend/kernels/elementwise_functions/erf.hpp
index 2ec7960bea97..a292a0a3f3f6 100644
--- a/dpnp/backend/kernels/elementwise_functions/erf.hpp
+++ b/dpnp/backend/kernels/elementwise_functions/erf.hpp
@@ -29,30 +29,46 @@
 
 #include <sycl/sycl.hpp>
 
-namespace dpnp::kernels::erf
+namespace dpnp::kernels::erfs
 {
-template <typename argT, typename Tp>
-struct ErfFunctor
+template <typename OpT, typename ArgT, typename ResT>
+struct BaseFunctor
 {
-    // is function constant for given argT
+    // is function constant for given ArgT
     using is_constant = typename std::false_type;
     // constant value, if constant
-    // constexpr Tp constant_value = Tp{};
+    // constexpr ResT constant_value = ResT{};
     // is function defined for sycl::vec
     using supports_vec = typename std::false_type;
-    // do both argT and Tp support subgroup store/load operation
+    // do both ArgT and ResT support subgroup store/load operation
     using supports_sg_loadstore = typename std::true_type;
 
-    Tp operator()(const argT &x) const
+    ResT operator()(const ArgT &x) const
     {
-        if constexpr (std::is_same_v<argT, sycl::half> &&
-                      std::is_same_v<Tp, float>) {
+        if constexpr (std::is_same_v<ArgT, sycl::half> &&
+                      std::is_same_v<ResT, float>) {
             // cast sycl::half to float for accuracy reasons
-            return sycl::erf(float(x));
+            return OpT::apply(float(x));
         }
         else {
-            return sycl::erf(x);
+            return OpT::apply(x);
         }
     }
 };
-} // namespace dpnp::kernels::erf
+
+#define MACRO_DEFINE_FUNCTOR(__name__, __f_name__)                             \
+    struct __f_name__##Op                                                      \
+    {                                                                          \
+        template <typename Tp>                                                 \
+        static Tp apply(const Tp &x)                                           \
+        {                                                                      \
+            return sycl::__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);
+} // namespace dpnp::kernels::erfs
diff --git a/dpnp/backend/src/dpnp_utils.hpp b/dpnp/backend/src/dpnp_utils.hpp
index d3a0d4659c70..7111b212b204 100644
--- a/dpnp/backend/src/dpnp_utils.hpp
+++ b/dpnp/backend/src/dpnp_utils.hpp
@@ -53,16 +53,6 @@
 #define __INTEL_MKL_2023_0_0_VERSION_REQUIRED 20230000
 #endif
 
-/**
- * Version of Intel MKL at which transition to OneMKL release 2023.2.0 occurs.
- *
- * @note with OneMKL=2023.1.0 the call of oneapi::mkl::vm::div() was dead
- * locked inside ~usm_wrapper_to_host()->{...; q_->wait_and_throw(); ...}
- */
-#ifndef __INTEL_MKL_2023_2_0_VERSION_REQUIRED
-#define __INTEL_MKL_2023_2_0_VERSION_REQUIRED 20230002L
-#endif
-
 /**
  * @defgroup BACKEND_UTILS Backend C++ library utilities
  * @{
diff --git a/dpnp/special/__init__.py b/dpnp/special/__init__.py
index 484fe38c29ef..8e1a3fcf24e0 100644
--- a/dpnp/special/__init__.py
+++ b/dpnp/special/__init__.py
@@ -41,8 +41,10 @@
 
 from ._erf import (
     erf,
+    erfc,
 )
 
 __all__ = [
     "erf",
+    "erfc",
 ]
diff --git a/dpnp/special/_erf.py b/dpnp/special/_erf.py
index 4629d3bbf8e2..6060f09e943b 100644
--- a/dpnp/special/_erf.py
+++ b/dpnp/special/_erf.py
@@ -42,7 +42,7 @@
 import dpnp.backend.extensions.ufunc._ufunc_impl as ufi
 from dpnp.dpnp_algo.dpnp_elementwise_common import DPNPUnaryFunc
 
-__all__ = ["erf"]
+__all__ = ["erf", "erfc"]
 
 
 # pylint: disable=too-few-public-methods
@@ -85,6 +85,11 @@ def __call__(self, x, out=None):  # pylint: disable=signature-differs
 out : {dpnp.ndarray, usm_ndarray}, optional
     Optional output array for the function values.
 
+Returns
+-------
+out : dpnp.ndarray
+    The values of the error function at the given points `x`.
+
 See Also
 --------
 :obj:`dpnp.special.erfc` : Complementary error function.
@@ -121,3 +126,48 @@ def __call__(self, x, out=None):  # pylint: disable=signature-differs
     mkl_fn_to_call="_mkl_erf_to_call",
     mkl_impl_fn="_erf",
 )
+
+_ERFC_DOCSTRING = r"""
+Calculates the complementary error function of a given input array.
+
+It is defined as :math:`1 - \operatorname{erf}(x)`.
+
+For full documentation refer to :obj:`scipy.special.erfc`.
+
+Parameters
+----------
+x : {dpnp.ndarray, usm_ndarray}
+    Input array, expected to have a real-valued floating-point data type.
+out : {dpnp.ndarray, usm_ndarray}, optional
+    Optional output array for the function values.
+
+Returns
+-------
+out : dpnp.ndarray
+    The values of the complementary error function at the given points `x`.
+
+See Also
+--------
+:obj:`dpnp.special.erf` : Gauss error function.
+:obj:`dpnp.special.erfinv` : Inverse of the error function.
+:obj:`dpnp.special.erfcinv` : Inverse of the complementary error function.
+:obj:`dpnp.special.erfcx` : Scaled complementary error function.
+:obj:`dpnp.special.erfi` : Imaginary error function.
+
+Examples
+--------
+>>> import dpnp as np
+>>> x = np.linspace(-3, 3, num=5)
+>>> np.special.erfc(x)
+array([[-0.99997791, -0.96610515,  0.        ,  0.96610515,  0.99997791])
+
+"""
+
+erfc = DPNPErf(
+    "erfc",
+    ufi._erf_result_type,
+    ufi._erfc,
+    _ERFC_DOCSTRING,
+    mkl_fn_to_call="_mkl_erf_to_call",
+    mkl_impl_fn="_erfc",
+)
diff --git a/dpnp/tests/test_special.py b/dpnp/tests/test_special.py
index 4dd0b72a7e8b..b7eb159b127d 100644
--- a/dpnp/tests/test_special.py
+++ b/dpnp/tests/test_special.py
@@ -14,18 +14,19 @@
 
 
 @with_requires("scipy")
-class TestErf:
+@pytest.mark.parametrize("func", ["erf", "erfc"])
+class TestCommon:
     @pytest.mark.parametrize(
         "dt", get_all_dtypes(no_none=True, no_float16=False, no_complex=True)
     )
-    def test_basic(self, dt):
+    def test_basic(self, func, dt):
         import scipy.special
 
         a = generate_random_numpy_array((2, 5), dtype=dt)
         ia = dpnp.array(a)
 
-        result = dpnp.special.erf(ia)
-        expected = scipy.special.erf(a)
+        result = getattr(dpnp.special, func)(ia)
+        expected = getattr(scipy.special, func)(a)
 
         # scipy >= 0.16.0 returns float64, but dpnp returns float32
         to_float32 = dt in (dpnp.bool, dpnp.float16)
@@ -34,29 +35,48 @@ def test_basic(self, dt):
             result, expected, check_only_type_kind=only_type_kind
         )
 
-    def test_nan_inf(self):
+    def test_nan_inf(self, func):
         import scipy.special
 
         a = numpy.array([numpy.nan, -numpy.inf, numpy.inf])
         ia = dpnp.array(a)
 
-        result = dpnp.special.erf(ia)
-        expected = scipy.special.erf(a)
+        result = getattr(dpnp.special, func)(ia)
+        expected = getattr(scipy.special, func)(a)
         assert_allclose(result, expected)
 
-    def test_zeros(self):
+    def test_zeros(self, func):
         import scipy.special
 
         a = numpy.array([0.0, -0.0])
         ia = dpnp.array(a)
 
-        result = dpnp.special.erf(ia)
-        expected = scipy.special.erf(a)
+        result = getattr(dpnp.special, func)(ia)
+        expected = getattr(scipy.special, func)(a)
         assert_allclose(result, expected)
         assert_equal(dpnp.signbit(result), numpy.signbit(expected))
 
     @pytest.mark.parametrize("dt", get_complex_dtypes())
-    def test_complex(self, dt):
+    def test_complex(self, func, dt):
         x = dpnp.empty(5, dtype=dt)
         with pytest.raises(ValueError):
-            dpnp.special.erf(x)
+            getattr(dpnp.special, func)(x)
+
+
+class TestConsistency:
+
+    def test_erfc(self):
+        # TODO: replace with dpnp.random.RandomState, once pareto is added
+        rng = numpy.random.RandomState(1234)
+        n = 10000
+        a = rng.pareto(0.02, n) * (2 * rng.randint(0, 2, n) - 1)
+        a = dpnp.array(a)
+
+        res = 1 - dpnp.special.erf(a)
+        mask = dpnp.isfinite(res)
+        a = a[mask]
+
+        tol = 8 * dpnp.finfo(a).resolution
+        assert dpnp.allclose(
+            dpnp.special.erfc(a), res[mask], rtol=tol, atol=tol
+        )
diff --git a/dpnp/tests/test_strides.py b/dpnp/tests/test_strides.py
index e124b0b7c162..607934c472ef 100644
--- a/dpnp/tests/test_strides.py
+++ b/dpnp/tests/test_strides.py
@@ -17,7 +17,7 @@
     get_integer_float_dtypes,
     numpy_version,
 )
-from .third_party.cupy.testing import installed, with_requires
+from .third_party.cupy.testing import with_requires
 
 
 @pytest.mark.usefixtures("suppress_divide_invalid_numpy_warnings")
@@ -167,20 +167,17 @@ def test_reduce_hypot(dtype, stride):
 
 
 @with_requires("scipy")
-@pytest.mark.parametrize("dtype", get_float_dtypes(no_float16=False))
+@pytest.mark.parametrize("func", ["erf", "erfc"])
 @pytest.mark.parametrize("stride", [2, -1, -3])
-def test_erf(dtype, stride):
+def test_erf_funcs(func, stride):
     import scipy.special
 
-    x = generate_random_numpy_array(10, dtype=dtype)
+    x = generate_random_numpy_array(10)
     a, ia = x[::stride], dpnp.array(x)[::stride]
 
-    result = dpnp.special.erf(ia)
-    expected = scipy.special.erf(a)
-
-    # scipy >= 0.16.0 returns float64, but dpnp returns float32
-    only_type_kind = installed("scipy>=0.16.0") and (dtype == dpnp.float16)
-    assert_dtype_allclose(result, expected, check_only_type_kind=only_type_kind)
+    result = getattr(dpnp.special, func)(ia)
+    expected = getattr(scipy.special, func)(a)
+    assert_dtype_allclose(result, expected)
 
 
 @pytest.mark.filterwarnings("ignore::RuntimeWarning")
diff --git a/dpnp/tests/test_sycl_queue.py b/dpnp/tests/test_sycl_queue.py
index f4299325762a..1ece65860d50 100644
--- a/dpnp/tests/test_sycl_queue.py
+++ b/dpnp/tests/test_sycl_queue.py
@@ -1489,11 +1489,12 @@ def test_interp(device, left, right, period):
     assert_sycl_queue_equal(result.sycl_queue, x.sycl_queue)
 
 
+@pytest.mark.parametrize("func", ["erf", "erfc"])
 @pytest.mark.parametrize("device", valid_dev, ids=dev_ids)
-def test_erf(device):
+def test_erf_funcs(func, device):
     x = dpnp.linspace(-3, 3, num=5, device=device)
 
-    result = dpnp.special.erf(x)
+    result = getattr(dpnp.special, func)(x)
     assert_sycl_queue_equal(result.sycl_queue, x.sycl_queue)
 
 
diff --git a/dpnp/tests/test_usm_type.py b/dpnp/tests/test_usm_type.py
index c17526649ab3..e97945e111a2 100644
--- a/dpnp/tests/test_usm_type.py
+++ b/dpnp/tests/test_usm_type.py
@@ -1296,10 +1296,11 @@ def test_choose(usm_type_x, usm_type_ind):
     assert z.usm_type == du.get_coerced_usm_type([usm_type_x, usm_type_ind])
 
 
+@pytest.mark.parametrize("func", ["erf", "erfc"])
 @pytest.mark.parametrize("usm_type", list_of_usm_types)
-def test_erf(usm_type):
+def test_erf_funcs(func, usm_type):
     x = dpnp.linspace(-3, 3, num=5, usm_type=usm_type)
-    y = dpnp.special.erf(x)
+    y = getattr(dpnp.special, func)(x)
     assert x.usm_type == y.usm_type == usm_type
 
 
diff --git a/dpnp/tests/third_party/cupyx/scipy_tests/special_tests/test_erf.py b/dpnp/tests/third_party/cupyx/scipy_tests/special_tests/test_erf.py
index 8ce8e6692987..c4a640a97d9e 100644
--- a/dpnp/tests/third_party/cupyx/scipy_tests/special_tests/test_erf.py
+++ b/dpnp/tests/third_party/cupyx/scipy_tests/special_tests/test_erf.py
@@ -19,11 +19,10 @@ def _boundary_inputs(boundary, rtol, atol):
 @testing.with_requires("scipy")
 class _TestBase:
 
-    # @testing.with_requires('scipy>=1.16.0')
+    @testing.with_requires("scipy>=1.16.0")
     def test_erf(self):
         self.check_unary("erf")
 
-    @pytest.mark.skip("erfc() is not supported yet")
     @testing.with_requires("scipy>=1.16.0")
     def test_erfc(self):
         self.check_unary("erfc")