diff --git a/dpnp/backend/extensions/vm/CMakeLists.txt b/dpnp/backend/extensions/vm/CMakeLists.txt
index 159ca57993ce..77f68d35ae83 100644
--- a/dpnp/backend/extensions/vm/CMakeLists.txt
+++ b/dpnp/backend/extensions/vm/CMakeLists.txt
@@ -52,6 +52,7 @@ set(_elementwise_sources
     ${CMAKE_CURRENT_SOURCE_DIR}/log1p.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/log2.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/mul.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/nextafter.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/pow.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/rint.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/sin.cpp
diff --git a/dpnp/backend/extensions/vm/abs.cpp b/dpnp/backend/extensions/vm/abs.cpp
index 7eb7086de85e..6857c2cdbe7d 100644
--- a/dpnp/backend/extensions/vm/abs.cpp
+++ b/dpnp/backend/extensions/vm/abs.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/acos.cpp b/dpnp/backend/extensions/vm/acos.cpp
index ab744bf99c44..fff7708c5f30 100644
--- a/dpnp/backend/extensions/vm/acos.cpp
+++ b/dpnp/backend/extensions/vm/acos.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/acosh.cpp b/dpnp/backend/extensions/vm/acosh.cpp
index 2cab39313d20..0b7c91f2a670 100644
--- a/dpnp/backend/extensions/vm/acosh.cpp
+++ b/dpnp/backend/extensions/vm/acosh.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/add.cpp b/dpnp/backend/extensions/vm/add.cpp
index c174bf73a99d..5b17da61b3fa 100644
--- a/dpnp/backend/extensions/vm/add.cpp
+++ b/dpnp/backend/extensions/vm/add.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/asin.cpp b/dpnp/backend/extensions/vm/asin.cpp
index afbb868e8cca..522bec92ede7 100644
--- a/dpnp/backend/extensions/vm/asin.cpp
+++ b/dpnp/backend/extensions/vm/asin.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/asinh.cpp b/dpnp/backend/extensions/vm/asinh.cpp
index 0f70c3cb5010..294f30589fd7 100644
--- a/dpnp/backend/extensions/vm/asinh.cpp
+++ b/dpnp/backend/extensions/vm/asinh.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/atan.cpp b/dpnp/backend/extensions/vm/atan.cpp
index 59f7064ef156..a6914d3d8a4c 100644
--- a/dpnp/backend/extensions/vm/atan.cpp
+++ b/dpnp/backend/extensions/vm/atan.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/atan2.cpp b/dpnp/backend/extensions/vm/atan2.cpp
index 4820a9623f0c..41ad849e7f4d 100644
--- a/dpnp/backend/extensions/vm/atan2.cpp
+++ b/dpnp/backend/extensions/vm/atan2.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/atanh.cpp b/dpnp/backend/extensions/vm/atanh.cpp
index bd32d25f2a6b..b274bb32b7c9 100644
--- a/dpnp/backend/extensions/vm/atanh.cpp
+++ b/dpnp/backend/extensions/vm/atanh.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/cbrt.cpp b/dpnp/backend/extensions/vm/cbrt.cpp
index 88bc82824180..25e5955dd2e2 100644
--- a/dpnp/backend/extensions/vm/cbrt.cpp
+++ b/dpnp/backend/extensions/vm/cbrt.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/ceil.cpp b/dpnp/backend/extensions/vm/ceil.cpp
index 14e7234a54c9..777ed1d03cdc 100644
--- a/dpnp/backend/extensions/vm/ceil.cpp
+++ b/dpnp/backend/extensions/vm/ceil.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/conj.cpp b/dpnp/backend/extensions/vm/conj.cpp
index edfb4384dad0..885a3fe1a397 100644
--- a/dpnp/backend/extensions/vm/conj.cpp
+++ b/dpnp/backend/extensions/vm/conj.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/cos.cpp b/dpnp/backend/extensions/vm/cos.cpp
index e7925cc32987..967a0757f913 100644
--- a/dpnp/backend/extensions/vm/cos.cpp
+++ b/dpnp/backend/extensions/vm/cos.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/cosh.cpp b/dpnp/backend/extensions/vm/cosh.cpp
index bb883c97c33e..40989b176cbb 100644
--- a/dpnp/backend/extensions/vm/cosh.cpp
+++ b/dpnp/backend/extensions/vm/cosh.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/div.cpp b/dpnp/backend/extensions/vm/div.cpp
index 5fb7122a76c2..2686456ae55a 100644
--- a/dpnp/backend/extensions/vm/div.cpp
+++ b/dpnp/backend/extensions/vm/div.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/exp.cpp b/dpnp/backend/extensions/vm/exp.cpp
index b7f8d4422d18..5bf091f18b75 100644
--- a/dpnp/backend/extensions/vm/exp.cpp
+++ b/dpnp/backend/extensions/vm/exp.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/exp2.cpp b/dpnp/backend/extensions/vm/exp2.cpp
index 8b5d7a7c5ff3..df6d06fd22a4 100644
--- a/dpnp/backend/extensions/vm/exp2.cpp
+++ b/dpnp/backend/extensions/vm/exp2.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/expm1.cpp b/dpnp/backend/extensions/vm/expm1.cpp
index b27668ba7c48..5409fd06270c 100644
--- a/dpnp/backend/extensions/vm/expm1.cpp
+++ b/dpnp/backend/extensions/vm/expm1.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/floor.cpp b/dpnp/backend/extensions/vm/floor.cpp
index 8a32f40e0ffb..f56c86949244 100644
--- a/dpnp/backend/extensions/vm/floor.cpp
+++ b/dpnp/backend/extensions/vm/floor.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/fmax.cpp b/dpnp/backend/extensions/vm/fmax.cpp
index b711516f6797..5899c4d6d70f 100644
--- a/dpnp/backend/extensions/vm/fmax.cpp
+++ b/dpnp/backend/extensions/vm/fmax.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/fmin.cpp b/dpnp/backend/extensions/vm/fmin.cpp
index 3b288216c921..b0817ea9c769 100644
--- a/dpnp/backend/extensions/vm/fmin.cpp
+++ b/dpnp/backend/extensions/vm/fmin.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/fmod.cpp b/dpnp/backend/extensions/vm/fmod.cpp
index e985492de047..bdc873c4658a 100644
--- a/dpnp/backend/extensions/vm/fmod.cpp
+++ b/dpnp/backend/extensions/vm/fmod.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/hypot.cpp b/dpnp/backend/extensions/vm/hypot.cpp
index 50ca178c37c3..217ba09f162d 100644
--- a/dpnp/backend/extensions/vm/hypot.cpp
+++ b/dpnp/backend/extensions/vm/hypot.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/ln.cpp b/dpnp/backend/extensions/vm/ln.cpp
index 2eb321a3777a..ed06faf340fb 100644
--- a/dpnp/backend/extensions/vm/ln.cpp
+++ b/dpnp/backend/extensions/vm/ln.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/log10.cpp b/dpnp/backend/extensions/vm/log10.cpp
index e685e5fce601..fe85ef9c0df4 100644
--- a/dpnp/backend/extensions/vm/log10.cpp
+++ b/dpnp/backend/extensions/vm/log10.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/log1p.cpp b/dpnp/backend/extensions/vm/log1p.cpp
index 2db1491e5ebd..b88442679c4e 100644
--- a/dpnp/backend/extensions/vm/log1p.cpp
+++ b/dpnp/backend/extensions/vm/log1p.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/log2.cpp b/dpnp/backend/extensions/vm/log2.cpp
index a6800185c256..4a176dae2a43 100644
--- a/dpnp/backend/extensions/vm/log2.cpp
+++ b/dpnp/backend/extensions/vm/log2.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/mul.cpp b/dpnp/backend/extensions/vm/mul.cpp
index de59d087f516..e35d41849588 100644
--- a/dpnp/backend/extensions/vm/mul.cpp
+++ b/dpnp/backend/extensions/vm/mul.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/nextafter.cpp b/dpnp/backend/extensions/vm/nextafter.cpp
new file mode 100644
index 000000000000..02c863ecc703
--- /dev/null
+++ b/dpnp/backend/extensions/vm/nextafter.cpp
@@ -0,0 +1,166 @@
+//*****************************************************************************
+// Copyright (c) 2024, 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 <oneapi/mkl.hpp>
+#include <sycl/sycl.hpp>
+
+#include "dpctl4pybind11.hpp"
+
+#include "common.hpp"
+#include "nextafter.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::nextafter<T> function.
+ *
+ * @tparam T Type of input vectors `a` and `b` and of result vector `y`.
+ */
+template <typename T1, typename T2>
+struct OutputType
+{
+    using value_type = typename std::disjunction<
+        td_ns::BinaryTypeMapResultEntry<T1, double, T2, double, double>,
+        td_ns::BinaryTypeMapResultEntry<T1, float, T2, float, float>,
+        td_ns::DefaultResultEntry<void>>::result_type;
+};
+
+template <typename T1, typename T2>
+static sycl::event
+    nextafter_contig_impl(sycl::queue &exec_q,
+                          std::size_t in_n,
+                          const char *in_a,
+                          py::ssize_t a_offset,
+                          const char *in_b,
+                          py::ssize_t b_offset,
+                          char *out_y,
+                          py::ssize_t out_offset,
+                          const std::vector<sycl::event> &depends)
+{
+    tu_ns::validate_type_for_device<T1>(exec_q);
+    tu_ns::validate_type_for_device<T2>(exec_q);
+
+    if ((a_offset != 0) || (b_offset != 0) || (out_offset != 0)) {
+        throw std::runtime_error("Arrays offsets have to be equals to 0");
+    }
+
+    std::int64_t n = static_cast<std::int64_t>(in_n);
+    const T1 *a = reinterpret_cast<const T1 *>(in_a);
+    const T2 *b = reinterpret_cast<const T2 *>(in_b);
+
+    using resTy = typename OutputType<T1, T2>::value_type;
+    resTy *y = reinterpret_cast<resTy *>(out_y);
+
+    return mkl_vm::nextafter(
+        exec_q,
+        n, // number of elements to be calculated
+        a, // pointer `a` containing 1st input vector of size n
+        b, // pointer `b` containing 2nd input vector of size n
+        y, // pointer `y` to the output vector of size n
+        depends);
+}
+
+using ew_cmn_ns::binary_contig_impl_fn_ptr_t;
+using ew_cmn_ns::binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t;
+using ew_cmn_ns::binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t;
+using ew_cmn_ns::binary_strided_impl_fn_ptr_t;
+
+static int output_typeid_vector[td_ns::num_types][td_ns::num_types];
+static binary_contig_impl_fn_ptr_t contig_dispatch_vector[td_ns::num_types]
+                                                         [td_ns::num_types];
+
+MACRO_POPULATE_DISPATCH_TABLES(nextafter);
+} // namespace impl
+
+void init_nextafter(py::module_ m)
+{
+    using arrayT = dpctl::tensor::usm_ndarray;
+    using event_vecT = std::vector<sycl::event>;
+
+    impl::populate_dispatch_tables();
+    using impl::contig_dispatch_vector;
+    using impl::output_typeid_vector;
+
+    auto nextafter_pyapi = [&](sycl::queue &exec_q, const arrayT &src1,
+                               const arrayT &src2, const arrayT &dst,
+                               const event_vecT &depends = {}) {
+        return py_int::py_binary_ufunc(
+            src1, src2, dst, exec_q, depends, output_typeid_vector,
+            contig_dispatch_vector,
+            // no support of strided implementation in OneMKL
+            td_ns::NullPtrTable<impl::binary_strided_impl_fn_ptr_t>{},
+            // no support of C-contig row with broadcasting in OneMKL
+            td_ns::NullPtrTable<
+                impl::
+                    binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{},
+            td_ns::NullPtrTable<
+                impl::
+                    binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{});
+    };
+    m.def(
+        "_nextafter", nextafter_pyapi,
+        "Call `nextafter` function from OneMKL VM library to return `dst` of "
+        "elements containing the next representable floating-point values "
+        "following the values from the elements of `src1` in the direction of "
+        "the corresponding elements of `src2`",
+        py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"), py::arg("dst"),
+        py::arg("depends") = py::list());
+
+    auto nextafter_need_to_call_pyapi = [&](sycl::queue &exec_q,
+                                            const arrayT &src1,
+                                            const arrayT &src2,
+                                            const arrayT &dst) {
+        return py_internal::need_to_call_binary_ufunc(exec_q, src1, src2, dst,
+                                                      output_typeid_vector,
+                                                      contig_dispatch_vector);
+    };
+    m.def("_mkl_nextafter_to_call", nextafter_need_to_call_pyapi,
+          "Check input arguments to answer if `nextafter` function from "
+          "OneMKL VM library can be used",
+          py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"),
+          py::arg("dst"));
+}
+} // namespace dpnp::extensions::vm
diff --git a/dpnp/backend/extensions/vm/nextafter.hpp b/dpnp/backend/extensions/vm/nextafter.hpp
new file mode 100644
index 000000000000..0b968764ec62
--- /dev/null
+++ b/dpnp/backend/extensions/vm/nextafter.hpp
@@ -0,0 +1,35 @@
+//*****************************************************************************
+// Copyright (c) 2024, Intel Corporation
+// All rights reserved.
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+// - Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// - Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
+// THE POSSIBILITY OF SUCH DAMAGE.
+//*****************************************************************************
+
+#pragma once
+
+#include <pybind11/pybind11.h>
+
+namespace py = pybind11;
+
+namespace dpnp::extensions::vm
+{
+void init_nextafter(py::module_ m);
+} // namespace dpnp::extensions::vm
diff --git a/dpnp/backend/extensions/vm/pow.cpp b/dpnp/backend/extensions/vm/pow.cpp
index 491b86f79469..7d2b2ea1000a 100644
--- a/dpnp/backend/extensions/vm/pow.cpp
+++ b/dpnp/backend/extensions/vm/pow.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/rint.cpp b/dpnp/backend/extensions/vm/rint.cpp
index ee0edbecd23e..38c3e0b32443 100644
--- a/dpnp/backend/extensions/vm/rint.cpp
+++ b/dpnp/backend/extensions/vm/rint.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/sin.cpp b/dpnp/backend/extensions/vm/sin.cpp
index 55d9f8ed301e..a1401961ca1c 100644
--- a/dpnp/backend/extensions/vm/sin.cpp
+++ b/dpnp/backend/extensions/vm/sin.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/sinh.cpp b/dpnp/backend/extensions/vm/sinh.cpp
index f8ddbc580ebc..d71b94b7cabb 100644
--- a/dpnp/backend/extensions/vm/sinh.cpp
+++ b/dpnp/backend/extensions/vm/sinh.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/sqr.cpp b/dpnp/backend/extensions/vm/sqr.cpp
index f42427ea00fc..8451c2911a79 100644
--- a/dpnp/backend/extensions/vm/sqr.cpp
+++ b/dpnp/backend/extensions/vm/sqr.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/sqrt.cpp b/dpnp/backend/extensions/vm/sqrt.cpp
index 70ebbf298fd3..720285f660f6 100644
--- a/dpnp/backend/extensions/vm/sqrt.cpp
+++ b/dpnp/backend/extensions/vm/sqrt.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/sub.cpp b/dpnp/backend/extensions/vm/sub.cpp
index 8bfc477bfa79..232523e40b54 100644
--- a/dpnp/backend/extensions/vm/sub.cpp
+++ b/dpnp/backend/extensions/vm/sub.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/tan.cpp b/dpnp/backend/extensions/vm/tan.cpp
index 250c38387227..65fa712c6cb3 100644
--- a/dpnp/backend/extensions/vm/tan.cpp
+++ b/dpnp/backend/extensions/vm/tan.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/tanh.cpp b/dpnp/backend/extensions/vm/tanh.cpp
index d0e9ecc1669a..94a2d9303015 100644
--- a/dpnp/backend/extensions/vm/tanh.cpp
+++ b/dpnp/backend/extensions/vm/tanh.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/trunc.cpp b/dpnp/backend/extensions/vm/trunc.cpp
index f47da825719c..54c373a3c1db 100644
--- a/dpnp/backend/extensions/vm/trunc.cpp
+++ b/dpnp/backend/extensions/vm/trunc.cpp
@@ -43,16 +43,15 @@
 
 namespace dpnp::extensions::vm
 {
-namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
 namespace py = pybind11;
 namespace py_int = dpnp::extensions::py_internal;
 namespace td_ns = dpctl::tensor::type_dispatch;
-namespace tu_ns = dpctl::tensor::type_utils;
 
 namespace impl
 {
-// OneMKL namespace with VM functions
-namespace mkl_vm = oneapi::mkl::vm;
+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
diff --git a/dpnp/backend/extensions/vm/vm_py.cpp b/dpnp/backend/extensions/vm/vm_py.cpp
index 4491703957a7..72040c1632b7 100644
--- a/dpnp/backend/extensions/vm/vm_py.cpp
+++ b/dpnp/backend/extensions/vm/vm_py.cpp
@@ -55,6 +55,7 @@
 #include "log1p.hpp"
 #include "log2.hpp"
 #include "mul.hpp"
+#include "nextafter.hpp"
 #include "pow.hpp"
 #include "rint.hpp"
 #include "sin.hpp"
@@ -98,6 +99,7 @@ PYBIND11_MODULE(_vm_impl, m)
     vm_ns::init_log1p(m);
     vm_ns::init_log2(m);
     vm_ns::init_mul(m);
+    vm_ns::init_nextafter(m);
     vm_ns::init_pow(m);
     vm_ns::init_rint(m);
     vm_ns::init_sin(m);
diff --git a/dpnp/dpnp_iface_mathematical.py b/dpnp/dpnp_iface_mathematical.py
index ba84ada206c2..0ece8c0b5b84 100644
--- a/dpnp/dpnp_iface_mathematical.py
+++ b/dpnp/dpnp_iface_mathematical.py
@@ -111,6 +111,7 @@
     "modf",
     "multiply",
     "negative",
+    "nextafter",
     "positive",
     "power",
     "prod",
@@ -2359,6 +2360,64 @@ def modf(x1, **kwargs):
 )
 
 
+_NEXTAFTER_DOCSTRING = """
+Return the next floating-point value after `x1` towards `x2`, element-wise.
+
+For full documentation refer to :obj:`numpy.nextafter`.
+
+Parameters
+----------
+x1 : {dpnp.ndarray, usm_ndarray, scalar}
+    Values to find the next representable value of.
+    Both inputs `x1` and `x2` can not be scalars at the same time.
+x2 : {dpnp.ndarray, usm_ndarray, scalar}
+    The direction where to look for the next representable value of `x1`.
+    Both inputs `x1` and `x2` can not be scalars at the same time.
+out : {None, dpnp.ndarray, usm_ndarray}, optional
+    Output array to populate. Array must have the correct shape and
+    the expected data type.
+    Default: ``None``.
+order : {"C", "F", "A", "K"}, optional
+    Memory layout of the newly output array, if parameter `out` is ``None``.
+    Default: ``"K"``.
+
+Returns
+-------
+out : dpnp.ndarray
+    The next representable values of `x1` in the direction of `x2`. The data
+    type of the returned array is determined by the Type Promotion Rules.
+
+Limitations
+-----------
+Parameters `where` and `subok` are supported with their default values.
+Keyword argument `kwargs` is currently unsupported.
+Otherwise ``NotImplementedError`` exception will be raised.
+
+Examples
+--------
+>>> import dpnp as np
+>>> a = np.array(1, dtype=np.float32)
+>>> eps = np.finfo(a.dtype).eps
+>>> np.nextafter(a, 2) == eps + 1
+array(True)
+
+>>> a = np.array([1, 2], dtype=np.float32)
+>>> b = np.array([2, 1], dtype=np.float32)
+>>> c = np.array([eps + 1, 2 - eps])
+>>> np.nextafter(a, b) == c
+array([ True,  True])
+"""
+
+nextafter = DPNPBinaryFunc(
+    "nextafter",
+    ti._nextafter_result_type,
+    ti._nextafter,
+    _NEXTAFTER_DOCSTRING,
+    mkl_fn_to_call=vmi._mkl_nextafter_to_call,
+    mkl_impl_fn=vmi._nextafter,
+)
+
+
 _POSITIVE_DOCSTRING = """
 Computes the numerical positive for each element `x_i` of input array `x`.
 
@@ -2438,7 +2497,7 @@ def modf(x1, **kwargs):
     the expected data type.
     Default: ``None``.
 order : {"C", "F", "A", "K"}, optional
-    Output array, if parameter `out` is ``None``.
+    Memory layout of the newly output array, if parameter `out` is ``None``.
     Default: ``"K"``.
 
 Returns
diff --git a/tests/skipped_tests.tbl b/tests/skipped_tests.tbl
index 944a4bd122d4..ec4ebe8fe4e9 100644
--- a/tests/skipped_tests.tbl
+++ b/tests/skipped_tests.tbl
@@ -57,8 +57,6 @@ tests/test_umath.py::test_umaths[('ldexp', 'di')]
 tests/test_umath.py::test_umaths[('ldexp', 'dl')]
 tests/test_umath.py::test_umaths[('logaddexp2', 'ff')]
 tests/test_umath.py::test_umaths[('logaddexp2', 'dd')]
-tests/test_umath.py::test_umaths[('nextafter', 'ff')]
-tests/test_umath.py::test_umaths[('nextafter', 'dd')]
 tests/test_umath.py::test_umaths[('spacing', 'f')]
 tests/test_umath.py::test_umaths[('spacing', 'd')]
 
@@ -217,11 +215,6 @@ tests/third_party/cupy/manipulation_tests/test_dims.py::TestInvalidBroadcast_par
 tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_logaddexp2
 tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_logaddexp2_infinities
 
-tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_frexp
-tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_ldexp
-tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_nextafter_combination
-tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_nextafter_float
-
 tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_nan_to_num
 tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_nan_to_num_negative
 tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_nan_to_num_for_old_numpy
diff --git a/tests/skipped_tests_gpu.tbl b/tests/skipped_tests_gpu.tbl
index 61f981c2b9cc..24eabe21055e 100644
--- a/tests/skipped_tests_gpu.tbl
+++ b/tests/skipped_tests_gpu.tbl
@@ -37,8 +37,6 @@ tests/test_umath.py::test_umaths[('ldexp', 'di')]
 tests/test_umath.py::test_umaths[('ldexp', 'dl')]
 tests/test_umath.py::test_umaths[('logaddexp2', 'ff')]
 tests/test_umath.py::test_umaths[('logaddexp2', 'dd')]
-tests/test_umath.py::test_umaths[('nextafter', 'ff')]
-tests/test_umath.py::test_umaths[('nextafter', 'dd')]
 tests/test_umath.py::test_umaths[('spacing', 'f')]
 tests/test_umath.py::test_umaths[('spacing', 'd')]
 
@@ -268,11 +266,6 @@ tests/third_party/cupy/manipulation_tests/test_dims.py::TestInvalidBroadcast_par
 tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_logaddexp2
 tests/third_party/cupy/math_tests/test_explog.py::TestExplog::test_logaddexp2_infinities
 
-tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_frexp
-tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_ldexp
-tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_nextafter_combination
-tests/third_party/cupy/math_tests/test_floating.py::TestFloating::test_nextafter_float
-
 tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_nan_to_num
 tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_nan_to_num_negative
 tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_nan_to_num_for_old_numpy
diff --git a/tests/test_mathematical.py b/tests/test_mathematical.py
index 1fe62f5db1dc..94870a007fdd 100644
--- a/tests/test_mathematical.py
+++ b/tests/test_mathematical.py
@@ -1114,6 +1114,140 @@ def test_subtract(self, dtype, lhs, rhs):
         self._test_mathematical("subtract", dtype, lhs, rhs, check_type=False)
 
 
+class TestNextafter:
+    @pytest.mark.parametrize("dt", get_float_dtypes())
+    @pytest.mark.parametrize(
+        "val1, val2",
+        [
+            pytest.param(1, 2),
+            pytest.param(1, 0),
+            pytest.param(1, 1),
+        ],
+    )
+    def test_float(self, val1, val2, dt):
+        v1 = numpy.array(val1, dtype=dt)
+        v2 = numpy.array(val2, dtype=dt)
+        iv1, iv2 = dpnp.array(v1), dpnp.array(v2)
+
+        result = dpnp.nextafter(iv1, iv2)
+        expected = numpy.nextafter(v1, v2)
+        assert_equal(result, expected)
+
+    @pytest.mark.parametrize("dt", get_float_dtypes())
+    def test_float_nan(self, dt):
+        a = numpy.array(1, dtype=dt)
+        ia = dpnp.array(a)
+
+        result = dpnp.nextafter(ia, dpnp.nan)
+        expected = numpy.nextafter(a, numpy.nan)
+        assert_equal(result, expected)
+
+        result = dpnp.nextafter(dpnp.nan, ia)
+        expected = numpy.nextafter(numpy.nan, a)
+        assert_equal(result, expected)
+
+    @pytest.mark.parametrize("val", [0x7C00, 0x8000], ids=["val1", "val2"])
+    def test_f16_strides(self, val):
+        a = numpy.arange(val, dtype=numpy.uint16).astype(numpy.float16)
+        hinf = numpy.array((numpy.inf,), dtype=numpy.float16)
+        ia, ihinf = dpnp.array(a), dpnp.array(hinf)
+
+        result = dpnp.nextafter(ia[:-1], ihinf)
+        expected = numpy.nextafter(a[:-1], hinf)
+        assert_equal(result, expected)
+
+        result = dpnp.nextafter(ia[0], -ihinf)
+        expected = numpy.nextafter(a[0], -hinf)
+        assert_equal(result, expected)
+
+        result = dpnp.nextafter(ia[1:], -ihinf)
+        expected = numpy.nextafter(a[1:], -hinf)
+        assert_equal(result, expected)
+
+    @pytest.mark.parametrize("val", [0x7C00, 0x8000], ids=["val1", "val2"])
+    def test_f16_array_inf(self, val):
+        a = numpy.arange(val, dtype=numpy.uint16).astype(numpy.float16)
+        hinf = numpy.array((numpy.inf,), dtype=numpy.float16)
+        ia, ihinf = dpnp.array(a), dpnp.array(hinf)
+
+        result = dpnp.nextafter(ihinf, ia)
+        expected = numpy.nextafter(hinf, a)
+        assert_equal(result, expected)
+
+        result = dpnp.nextafter(-ihinf, ia)
+        expected = numpy.nextafter(-hinf, a)
+        assert_equal(result, expected)
+
+    @pytest.mark.parametrize(
+        "sign1, sign2",
+        [
+            pytest.param(1, 1),
+            pytest.param(1, -1),
+            pytest.param(-1, 1),
+            pytest.param(-1, -1),
+        ],
+    )
+    def test_f16_inf(self, sign1, sign2):
+        hinf1 = numpy.array((sign1 * numpy.inf,), dtype=numpy.float16)
+        hinf2 = numpy.array((sign2 * numpy.inf,), dtype=numpy.float16)
+        ihinf1, ihinf2 = dpnp.array(hinf1), dpnp.array(hinf2)
+
+        result = dpnp.nextafter(ihinf1, ihinf2)
+        expected = numpy.nextafter(hinf1, hinf2)
+        assert_equal(result, expected)
+
+    @pytest.mark.parametrize("val", [0x7C00, 0x8000], ids=["val1", "val2"])
+    def test_f16_array_nan(self, val):
+        a = numpy.arange(val, dtype=numpy.uint16).astype(numpy.float16)
+        nan = numpy.array((numpy.nan,), dtype=numpy.float16)
+        ia, inan = dpnp.array(a), dpnp.array(nan)
+
+        result = dpnp.nextafter(ia, inan)
+        expected = numpy.nextafter(a, nan)
+        assert_equal(result, expected)
+
+        result = dpnp.nextafter(inan, ia)
+        expected = numpy.nextafter(nan, a)
+        assert_equal(result, expected)
+
+    @pytest.mark.parametrize(
+        "val1, val2",
+        [
+            pytest.param(numpy.nan, numpy.nan),
+            pytest.param(numpy.inf, numpy.nan),
+            pytest.param(numpy.nan, numpy.inf),
+        ],
+    )
+    def test_f16_inf_nan(self, val1, val2):
+        v1 = numpy.array((val1,), dtype=numpy.float16)
+        v2 = numpy.array((val2,), dtype=numpy.float16)
+        iv1, iv2 = dpnp.array(v1), dpnp.array(v2)
+
+        result = dpnp.nextafter(iv1, iv2)
+        expected = numpy.nextafter(v1, v2)
+        assert_equal(result, expected)
+
+    @pytest.mark.parametrize(
+        "val, scalar",
+        [
+            pytest.param(65504, -numpy.inf),
+            pytest.param(-65504, numpy.inf),
+            pytest.param(numpy.inf, 0),
+            pytest.param(-numpy.inf, 0),
+            pytest.param(0, numpy.nan),
+            pytest.param(numpy.nan, 0),
+        ],
+    )
+    def test_f16_corner_values_with_scalar(self, val, scalar):
+        a = numpy.array(val, dtype=numpy.float16)
+        ia = dpnp.array(a)
+        scalar = numpy.float16(scalar)
+
+        result = dpnp.nextafter(ia, scalar)
+        expected = numpy.nextafter(a, scalar)
+        assert_equal(result, expected)
+
+
 @pytest.mark.usefixtures("suppress_divide_invalid_numpy_warnings")
 @pytest.mark.parametrize(
     "val_type", [bool, int, float], ids=["bool", "int", "float"]
diff --git a/tests/test_sycl_queue.py b/tests/test_sycl_queue.py
index 5b6287dc9f06..528d33084d0c 100644
--- a/tests/test_sycl_queue.py
+++ b/tests/test_sycl_queue.py
@@ -686,6 +686,7 @@ def test_reduce_hypot(device):
             [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0],
             [0.0, 1.0, 2.0, 0.0, 1.0, 2.0, 0.0, 1.0, 2.0],
         ),
+        pytest.param("nextafter", [1, 2], [2, 1]),
         pytest.param(
             "outer", [0.0, 1.0, 2.0, 3.0, 4.0, 5.0], [0.0, 1.0, 2.0, 0.0]
         ),
diff --git a/tests/test_usm_type.py b/tests/test_usm_type.py
index 6cc8d5edd39d..8abc0aaf3220 100644
--- a/tests/test_usm_type.py
+++ b/tests/test_usm_type.py
@@ -655,6 +655,7 @@ def test_1in_1out(func, data, usm_type):
         pytest.param("logaddexp", [[-1, 2, 5, 9]], [[4, -3, 2, -8]]),
         pytest.param("maximum", [0.0, 1.0, 2.0], [3.0, 4.0, 5.0]),
         pytest.param("minimum", [0.0, 1.0, 2.0], [3.0, 4.0, 5.0]),
+        pytest.param("nextafter", [1, 2], [2, 1]),
         pytest.param("searchsorted", [11, 12, 13, 14, 15], [-10, 20, 12, 13]),
         pytest.param(
             "tensordot",
diff --git a/tests/third_party/cupy/math_tests/test_floating.py b/tests/third_party/cupy/math_tests/test_floating.py
index 09531104a543..9b72537247bb 100644
--- a/tests/third_party/cupy/math_tests/test_floating.py
+++ b/tests/third_party/cupy/math_tests/test_floating.py
@@ -1,6 +1,7 @@
 import unittest
 
 import numpy
+import pytest
 
 import dpnp as cupy
 from tests.helper import has_support_aspect64
@@ -28,6 +29,7 @@ def test_copysign_float(self, xp, dtype):
         b = xp.array([-xp.inf, -3, -0.0, 0, 3, xp.inf], dtype=dtype)[None, :]
         return xp.copysign(a, b)
 
+    @pytest.mark.skip("ldexp() is not implemented yet")
     @testing.for_float_dtypes(name="ftype")
     @testing.for_dtypes(["i", "l"], name="itype")
     @testing.numpy_cupy_array_equal()
@@ -36,6 +38,7 @@ def test_ldexp(self, xp, ftype, itype):
         b = xp.array([-3, -2, -1, 0, 1, 2, 3], dtype=itype)
         return xp.ldexp(a, b)
 
+    @pytest.mark.skip("frexp() is not implemented yet")
     @testing.for_float_dtypes()
     def test_frexp(self, dtype):
         numpy_a = numpy.array(
@@ -50,7 +53,7 @@ def test_frexp(self, dtype):
         testing.assert_array_equal(cupy_c, numpy_c)
 
     @testing.for_all_dtypes_combination(("dtype_a", "dtype_b"), no_complex=True)
-    @testing.numpy_cupy_array_equal()
+    @testing.numpy_cupy_allclose(rtol=1e-06, type_check=has_support_aspect64())
     def test_nextafter_combination(self, xp, dtype_a, dtype_b):
         a = testing.shaped_arange((2, 3), xp, dtype_a)
         # skip 0 because cupy (may) handle denormals differently (-ftz=true)