Repository: IntelPython/numba-dpex (language: python)

Issue:
Queue inference for numba_dpex kernel args returns incorrect values
The following code should raise a `ComputeFollowsDataInferenceError`, however such an exception is not raised:

```python
import dpnp
import dpctl

import numba_dpex as dpex
from numba_dpex.core.exceptions import ComputeFollowsDataInferenceError


@dpex.kernel
def sum_kernel(a, b, c):
    i = dpex.get_global_id(0)
    c[i] = a[i] + b[i]

q1 = dpctl.SyclQueue()
q2 = dpctl.SyclQueue()

print(q1 == q2)
a = dpnp.ones(1, sycl_queue=q1)
b = dpnp.ones(1, sycl_queue=q2)
c = dpnp.empty_like(a)

sum_kernel[dpex.Range(1)](a, b, c)
```
Inside dispatcher when the type of the input arguments is inferred I see that all the args are inferred as `usm_ndarray` type with the same queue.


Discussion / hints:

Fix patch:
diff --git a/numba_dpex/__init__.py b/numba_dpex/__init__.py
index d043aea5ce..9ee692931d 100644
--- a/numba_dpex/__init__.py
+++ b/numba_dpex/__init__.py
@@ -94,7 +94,6 @@ def parse_sem_version(version_string: str) -> Tuple[int, int, int]:
 from numba import prange  # noqa E402
 
 import numba_dpex.core.dpjit_dispatcher  # noqa E402
-import numba_dpex.core.offload_dispatcher  # noqa E402
 
 # Initialize the _dpexrt_python extension
 import numba_dpex.core.runtime  # noqa E402
@@ -111,16 +110,11 @@ def parse_sem_version(version_string: str) -> Tuple[int, int, int]:
 # Re-export all type names
 from numba_dpex.core.types import *  # noqa E402
 from numba_dpex.dpnp_iface import dpnpimpl  # noqa E402
-from numba_dpex.retarget import offload_to_sycl_device  # noqa E402
 
 if config.HAS_NON_HOST_DEVICE:
     # Re export
     from .core.targets import dpjit_target, kernel_target
     from .decorators import dpjit, func, kernel
-
-    # We are importing dpnp stub module to make Numba recognize the
-    # module when we rename Numpy functions.
-    from .dpnp_iface.stubs import dpnp
     from .ocl.stubs import (
         GLOBAL_MEM_FENCE,
         LOCAL_MEM_FENCE,
@@ -150,4 +144,4 @@ def parse_sem_version(version_string: str) -> Tuple[int, int, int]:
 __version__ = get_versions()["version"]
 del get_versions
 
-__all__ = types.__all__ + ["offload_to_sycl_device"] + ["Range", "NdRange"]
+__all__ = types.__all__ + ["Range", "NdRange"]
diff --git a/numba_dpex/core/compiler.py b/numba_dpex/core/compiler.py
index cfb0e7d8ab..7c78ae9983 100644
--- a/numba_dpex/core/compiler.py
+++ b/numba_dpex/core/compiler.py
@@ -14,7 +14,6 @@
     UnreachableError,
 )
 from numba_dpex.core.pipelines.kernel_compiler import KernelCompiler
-from numba_dpex.core.pipelines.offload_compiler import OffloadCompiler
 
 
 @global_compiler_lock
@@ -73,10 +72,6 @@ def compile_with_dpex(
             pipeline_class=KernelCompiler,
         )
     elif isinstance(pyfunc, ir.FunctionIR):
-        # FIXME: Kernels in the form of Numba IR need to be compiled
-        # using the offload compiler due to them retaining parfor
-        # nodes due to the use of gufuncs. Once the kernel builder is
-        # ready we should be able to switch to the KernelCompiler.
         cres = compiler.compile_ir(
             typingctx=typingctx,
             targetctx=targetctx,
diff --git a/numba_dpex/core/offload_dispatcher.py b/numba_dpex/core/offload_dispatcher.py
deleted file mode 100644
index 43526612b4..0000000000
--- a/numba_dpex/core/offload_dispatcher.py
+++ /dev/null
@@ -1,60 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-from numba.core import compiler, dispatcher
-from numba.core.registry import cpu_target
-from numba.core.target_extension import dispatcher_registry, target_registry
-
-import numba_dpex.config as dpex_config
-from numba_dpex.core.targets.kernel_target import DPEX_KERNEL_TARGET_NAME
-
-
-class OffloadDispatcher(dispatcher.Dispatcher):
-    targetdescr = cpu_target
-
-    def __init__(
-        self,
-        py_func,
-        locals={},
-        targetoptions={},
-        impl_kind="direct",
-        pipeline_class=compiler.Compiler,
-    ):
-        if dpex_config.HAS_NON_HOST_DEVICE:
-            from numba_dpex.core.pipelines.offload_compiler import (
-                OffloadCompiler,
-            )
-
-            targetoptions["parallel"] = True
-            dispatcher.Dispatcher.__init__(
-                self,
-                py_func,
-                locals=locals,
-                targetoptions=targetoptions,
-                impl_kind=impl_kind,
-                pipeline_class=OffloadCompiler,
-            )
-        else:
-            print(
-                "--------------------------------------------------------------"
-            )
-            print(
-                "WARNING : DPEX pipeline ignored. Ensure drivers are installed."
-            )
-            print(
-                "--------------------------------------------------------------"
-            )
-            dispatcher.Dispatcher.__init__(
-                self,
-                py_func,
-                locals=locals,
-                targetoptions=targetoptions,
-                impl_kind=impl_kind,
-                pipeline_class=pipeline_class,
-            )
-
-
-dispatcher_registry[
-    target_registry[DPEX_KERNEL_TARGET_NAME]
-] = OffloadDispatcher
diff --git a/numba_dpex/core/passes/rename_numpy_functions_pass.py b/numba_dpex/core/passes/rename_numpy_functions_pass.py
deleted file mode 100644
index 5b93ca27d3..0000000000
--- a/numba_dpex/core/passes/rename_numpy_functions_pass.py
+++ /dev/null
@@ -1,361 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-from numba.core import ir, types
-from numba.core.compiler_machinery import FunctionPass, register_pass
-from numba.core.ir_utils import (
-    find_topo_order,
-    mk_unique_var,
-    remove_dead,
-    simplify_CFG,
-)
-
-import numba_dpex
-
-rewrite_function_name_map = {
-    # numpy
-    "all": (["numpy"], "all"),
-    "amax": (["numpy"], "amax"),
-    "amin": (["numpy"], "amin"),
-    "argmax": (["numpy"], "argmax"),
-    "argmin": (["numpy"], "argmin"),
-    "argsort": (["numpy"], "argsort"),
-    "cov": (["numpy"], "cov"),
-    "diagonal": (["numpy"], "diagonal"),
-    "max": (["numpy"], "max"),
-    "mean": (["numpy"], "mean"),
-    "median": (["numpy"], "median"),
-    "min": (["numpy"], "min"),
-    "partition": (["numpy"], "partition"),
-    "repeat": (["numpy"], "repeat"),
-    "trace": (["numpy"], "trace"),
-    "vdot": (["numpy"], "vdot"),
-    # random
-    "beta": (["random"], "beta"),
-    "binomial": (["random"], "binomial"),
-    "chisquare": (["random"], "chisquare"),
-    "exponential": (["random"], "exponential"),
-    "gamma": (["random"], "gamma"),
-    "geometric": (["random"], "geometric"),
-    "gumbel": (["random"], "gumbel"),
-    "hypergeometric": (["random"], "hypergeometric"),
-    "laplace": (["random"], "laplace"),
-    "lognormal": (["random"], "lognormal"),
-    "multinomial": (["random"], "multinomial"),
-    "multivariate_normal": (["random"], "multivariate_normal"),
-    "negative_binomial": (["random"], "negative_binomial"),
-    "normal": (["random"], "normal"),
-    "poisson": (["random"], "poisson"),
-    "rand": (["random"], "rand"),
-    "randint": (["random"], "randint"),
-    "random_integers": (["random"], "random_integers"),
-    "random_sample": (["random"], "random_sample"),
-    "random": (["random"], "random"),
-    "ranf": (["random"], "ranf"),
-    "rayleigh": (["random"], "rayleigh"),
-    "sample": (["random"], "sample"),
-    "standard_cauchy": (["random"], "standard_cauchy"),
-    "standard_exponential": (["random"], "standard_exponential"),
-    "standard_gamma": (["random"], "standard_gamma"),
-    "standard_normal": (["random"], "standard_normal"),
-    "uniform": (["random"], "uniform"),
-    "weibull": (["random"], "weibull"),
-    # linalg
-    "cholesky": (["linalg"], "cholesky"),
-    "det": (["linalg"], "det"),
-    "dot": (["numpy"], "dot"),
-    "eig": (["linalg"], "eig"),
-    "eigvals": (["linalg"], "eigvals"),
-    "matmul": (["numpy"], "matmul"),
-    "matrix_power": (["linalg"], "matrix_power"),
-    "matrix_rank": (["linalg"], "matrix_rank"),
-    "multi_dot": (["linalg"], "multi_dot"),
-    # transcendentals
-    "nanprod": (["numpy"], "nanprod"),
-    "nansum": (["numpy"], "nansum"),
-    "prod": (["numpy"], "prod"),
-    "sum": (["numpy"], "sum"),
-    # array ops
-    "copy": (["numpy"], "copy"),
-    "cumsum": (["numpy"], "cumsum"),
-    "cumprod": (["numpy"], "cumprod"),
-    "sort": (["numpy"], "sort"),
-    "take": (["numpy"], "take"),
-}
-
-
-class _RewriteNumPyOverloadedFunctionsImpl(object):
-    def __init__(
-        self, state, rewrite_function_name_map=rewrite_function_name_map
-    ):
-        self.state = state
-        self.function_name_map = rewrite_function_name_map
-
-    def run(self):
-        """
-        This function rewrites the name of NumPy functions that exist in self.function_name_map
-        e.g np.sum(a) would produce the following:
-
-        np.sum() --> numba_dpex.dpnp.sum()
-
-        ---------------------------------------------------------------------------------------
-        Numba IR Before Rewrite:
-        ---------------------------------------------------------------------------------------
-
-            $2load_global.0 = global(np: <module 'numpy' from 'numpy/__init__.py'>) ['$2load_global.0']
-            $4load_method.1 = getattr(value=$2load_global.0, attr=sum) ['$2load_global.0', '$4load_method.1']
-            $8call_method.3 = call $4load_method.1(a, func=$4load_method.1, args=[Var(a, test_rewrite.py:7)],
-                                                   kws=(), vararg=None) ['$4load_method.1', '$8call_method.3', 'a']
-
-        ---------------------------------------------------------------------------------------
-        Numba IR After Rewrite:
-        ---------------------------------------------------------------------------------------
-
-            $dpex_replaced_var.0 = global(numba_dpex: <module 'numba_dpex' from 'numba_dpex/__init__.py'>) ['$dpex_replaced_var.0']
-            $dpnp_var.1 = getattr(value=$dpex_replaced_var.0, attr=dpnp) ['$dpnp_var.1', '$dpex_replaced_var.0']
-            $4load_method.1 = getattr(value=$dpnp_var.1, attr=sum) ['$4load_method.1', '$dpnp_var.1']
-            $8call_method.3 = call $4load_method.1(a, func=$4load_method.1, args=[Var(a, test_rewrite.py:7)],
-                                                   kws=(), vararg=None) ['$4load_method.1', '$8call_method.3', 'a']
-
-        ---------------------------------------------------------------------------------------
-        """
-        func_ir = self.state.func_ir
-        blocks = func_ir.blocks
-        topo_order = find_topo_order(blocks)
-        replaced = False
-
-        for label in topo_order:
-            block = blocks[label]
-            saved_arr_arg = {}
-            new_body = []
-            for stmt in block.body:
-                if isinstance(stmt, ir.Assign) and isinstance(
-                    stmt.value, ir.Expr
-                ):
-                    lhs = stmt.target.name
-                    rhs = stmt.value
-                    # replace np.FOO with name from self.function_name_map["FOO"]
-                    # e.g. np.sum will be replaced with numba_dpex.dpnp.sum
-                    if (
-                        rhs.op == "getattr"
-                        and rhs.attr in self.function_name_map
-                    ):
-                        module_node = block.find_variable_assignment(
-                            rhs.value.name
-                        ).value
-                        if (
-                            isinstance(module_node, ir.Global)
-                            and module_node.value.__name__
-                            in self.function_name_map[rhs.attr][0]
-                        ) or (
-                            isinstance(module_node, ir.Expr)
-                            and module_node.attr
-                            in self.function_name_map[rhs.attr][0]
-                        ):
-                            rhs = stmt.value
-                            rhs.attr = self.function_name_map[rhs.attr][1]
-
-                            global_module = rhs.value
-                            saved_arr_arg[lhs] = global_module
-
-                            scope = global_module.scope
-                            loc = global_module.loc
-
-                            g_dpex_var = ir.Var(
-                                scope, mk_unique_var("$2load_global"), loc
-                            )
-                            # We are trying to rename np.function_name/np.linalg.function_name with
-                            # numba_dpex.dpnp.function_name.
-                            # Hence, we need to have a global variable representing module numba_dpex.
-                            # Next, we add attribute dpnp to global module numba_dpex to
-                            # represent numba_dpex.dpnp.
-                            g_dpex = ir.Global("numba_dpex", numba_dpex, loc)
-                            g_dpex_assign = ir.Assign(g_dpex, g_dpex_var, loc)
-
-                            dpnp_var = ir.Var(
-                                scope, mk_unique_var("$4load_attr"), loc
-                            )
-                            getattr_dpnp = ir.Expr.getattr(
-                                g_dpex_var, "dpnp", loc
-                            )
-                            dpnp_assign = ir.Assign(getattr_dpnp, dpnp_var, loc)
-
-                            rhs.value = dpnp_var
-                            new_body.append(g_dpex_assign)
-                            new_body.append(dpnp_assign)
-                            func_ir._definitions[dpnp_var.name] = [getattr_dpnp]
-                            func_ir._definitions[g_dpex_var.name] = [g_dpex]
-                            replaced = True
-
-                new_body.append(stmt)
-            block.body = new_body
-            return replaced
-
-
-@register_pass(mutates_CFG=True, analysis_only=False)
-class RewriteOverloadedNumPyFunctionsPass(FunctionPass):
-    _name = "dpex_rewrite_overloaded_functions_pass"
-
-    def __init__(self):
-        FunctionPass.__init__(self)
-
-        import numba_dpex.dpnp_iface.dpnp_array_ops_impl
-        import numba_dpex.dpnp_iface.dpnp_indexing
-        import numba_dpex.dpnp_iface.dpnp_linalgimpl
-        import numba_dpex.dpnp_iface.dpnp_logic
-        import numba_dpex.dpnp_iface.dpnp_manipulation
-        import numba_dpex.dpnp_iface.dpnp_randomimpl
-        import numba_dpex.dpnp_iface.dpnp_sort_search_countimpl
-        import numba_dpex.dpnp_iface.dpnp_statisticsimpl
-        import numba_dpex.dpnp_iface.dpnp_stubs_impl
-        import numba_dpex.dpnp_iface.dpnp_transcendentalsimpl
-        import numba_dpex.dpnp_iface.dpnpdecl
-
-    def run_pass(self, state):
-        rewrite_function_name_pass = _RewriteNumPyOverloadedFunctionsImpl(
-            state, rewrite_function_name_map
-        )
-
-        mutated = rewrite_function_name_pass.run()
-
-        if mutated:
-            remove_dead(
-                state.func_ir.blocks, state.func_ir.arg_names, state.func_ir
-            )
-        state.func_ir.blocks = simplify_CFG(state.func_ir.blocks)
-
-        return mutated
-
-
-def get_dpnp_func_typ(func):
-    from numba.core.typing.templates import builtin_registry
-
-    for k, v in builtin_registry.globals:
-        if k == func:
-            return v
-    raise RuntimeError("type for func ", func, " not found")
-
-
-class _RewriteNdarrayFunctionsImpl(object):
-    def __init__(
-        self, state, rewrite_function_name_map=rewrite_function_name_map
-    ):
-        self.state = state
-        self.function_name_map = rewrite_function_name_map
-        self.typemap = state.type_annotation.typemap
-        self.calltypes = state.type_annotation.calltypes
-
-    def run(self):
-        typingctx = self.state.typingctx
-
-        # save array arg to call
-        # call_varname -> array
-        func_ir = self.state.func_ir
-        blocks = func_ir.blocks
-        saved_arr_arg = {}
-        topo_order = find_topo_order(blocks)
-        replaced = False
-
-        for label in topo_order:
-            block = blocks[label]
-            new_body = []
-            for stmt in block.body:
-                if isinstance(stmt, ir.Assign) and isinstance(
-                    stmt.value, ir.Expr
-                ):
-                    lhs = stmt.target.name
-                    rhs = stmt.value
-                    # replace A.func with np.func, and save A in saved_arr_arg
-                    if (
-                        rhs.op == "getattr"
-                        and rhs.attr in self.function_name_map
-                        and isinstance(
-                            self.typemap[rhs.value.name], types.npytypes.Array
-                        )
-                    ):
-                        rhs = stmt.value
-                        arr = rhs.value
-                        saved_arr_arg[lhs] = arr
-                        scope = arr.scope
-                        loc = arr.loc
-
-                        g_dpex_var = ir.Var(
-                            scope, mk_unique_var("$load_global"), loc
-                        )
-                        self.typemap[g_dpex_var.name] = types.misc.Module(
-                            numba_dpex
-                        )
-                        g_dpex = ir.Global("numba_dpex", numba_dpex, loc)
-                        g_dpex_assign = ir.Assign(g_dpex, g_dpex_var, loc)
-
-                        dpnp_var = ir.Var(
-                            scope, mk_unique_var("$load_attr"), loc
-                        )
-                        self.typemap[dpnp_var.name] = types.misc.Module(
-                            numba_dpex.dpnp
-                        )
-                        getattr_dpnp = ir.Expr.getattr(g_dpex_var, "dpnp", loc)
-                        dpnp_assign = ir.Assign(getattr_dpnp, dpnp_var, loc)
-
-                        rhs.value = dpnp_var
-                        new_body.append(g_dpex_assign)
-                        new_body.append(dpnp_assign)
-
-                        func_ir._definitions[g_dpex_var.name] = [getattr_dpnp]
-                        func_ir._definitions[dpnp_var.name] = [getattr_dpnp]
-
-                        # update func var type
-                        func = getattr(numba_dpex.dpnp, rhs.attr)
-                        func_typ = get_dpnp_func_typ(func)
-
-                        self.typemap.pop(lhs)
-                        self.typemap[lhs] = func_typ
-                        replaced = True
-
-                    if rhs.op == "call" and rhs.func.name in saved_arr_arg:
-                        # add array as first arg
-                        arr = saved_arr_arg[rhs.func.name]
-                        # update call type signature to include array arg
-                        old_sig = self.calltypes.pop(rhs)
-                        # argsort requires kws for typing so sig.args can't be used
-                        # reusing sig.args since some types become Const in sig
-                        argtyps = old_sig.args[: len(rhs.args)]
-                        kwtyps = {
-                            name: self.typemap[v.name] for name, v in rhs.kws
-                        }
-                        self.calltypes[rhs] = self.typemap[
-                            rhs.func.name
-                        ].get_call_type(
-                            typingctx,
-                            [self.typemap[arr.name]] + list(argtyps),
-                            kwtyps,
-                        )
-                        rhs.args = [arr] + rhs.args
-
-                new_body.append(stmt)
-            block.body = new_body
-        return replaced
-
-
-@register_pass(mutates_CFG=True, analysis_only=False)
-class RewriteNdarrayFunctionsPass(FunctionPass):
-    _name = "dpex_rewrite_ndarray_functions_pass"
-
-    def __init__(self):
-        FunctionPass.__init__(self)
-
-    def run_pass(self, state):
-        rewrite_ndarray_function_name_pass = _RewriteNdarrayFunctionsImpl(
-            state, rewrite_function_name_map
-        )
-
-        mutated = rewrite_ndarray_function_name_pass.run()
-
-        if mutated:
-            remove_dead(
-                state.func_ir.blocks, state.func_ir.arg_names, state.func_ir
-            )
-        state.func_ir.blocks = simplify_CFG(state.func_ir.blocks)
-
-        return mutated
diff --git a/numba_dpex/core/pipelines/offload_compiler.py b/numba_dpex/core/pipelines/offload_compiler.py
deleted file mode 100644
index 017f26bf52..0000000000
--- a/numba_dpex/core/pipelines/offload_compiler.py
+++ /dev/null
@@ -1,204 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-from numba.core.compiler import CompilerBase
-from numba.core.compiler_machinery import PassManager
-from numba.core.typed_passes import (
-    AnnotateTypes,
-    InlineOverloads,
-    IRLegalization,
-    NopythonRewrites,
-    NoPythonSupportedFeatureValidation,
-    NopythonTypeInference,
-    PreLowerStripPhis,
-)
-from numba.core.untyped_passes import (
-    DeadBranchPrune,
-    FindLiterallyCalls,
-    FixupArgs,
-    GenericRewrites,
-    InlineClosureLikes,
-    InlineInlinables,
-    IRProcessing,
-    LiteralPropagationSubPipelinePass,
-    LiteralUnroll,
-    MakeFunctionToJitFunction,
-    ReconstructSSA,
-    RewriteSemanticConstants,
-    TranslateByteCode,
-    WithLifting,
-)
-
-from numba_dpex.core.exceptions import UnsupportedCompilationModeError
-from numba_dpex.core.passes.passes import (
-    DpexLowering,
-    DumpParforDiagnostics,
-    NoPythonBackend,
-    ParforPass,
-    PreParforPass,
-)
-from numba_dpex.core.passes.rename_numpy_functions_pass import (
-    RewriteNdarrayFunctionsPass,
-    RewriteOverloadedNumPyFunctionsPass,
-)
-from numba_dpex.parfor_diagnostics import ExtendedParforDiagnostics
-
-
-class _OffloadPassBuilder(object):
-    """
-    A pass builder for dpex's OffloAdCompiler that adds supports for
-    offloading parfors and other NumPy library calls to a SYCL device.
-
-    The pass builder does not implement pipelines for objectmode or interpreted
-    execution.
-    """
-
-    @staticmethod
-    def define_untyped_pipeline(state, name="dpex_offload_untyped"):
-        """Returns an untyped part of the nopython pipeline
-
-        The dpex offload pipeline's untyped passes are based on Numba's
-        compiler. Dpex addes extra passes to change specific numpy overloads
-        for which dpex provides alternate implementations.
-        """
-        pm = PassManager(name)
-        if state.func_ir is None:
-            pm.add_pass(TranslateByteCode, "analyzing bytecode")
-            pm.add_pass(FixupArgs, "fix up args")
-        pm.add_pass(IRProcessing, "processing IR")
-        pm.add_pass(WithLifting, "Handle with contexts")
-
-        # --- Begin dpex passes added to the untyped pipeline                --#
-
-        # The RewriteOverloadedNumPyFunctionsPass rewrites the module namespace
-        # of specific NumPy functions to dpnp, as we overload these functions
-        # differently.
-        pm.add_pass(
-            RewriteOverloadedNumPyFunctionsPass,
-            "Rewrite name of Numpy functions to overload already overloaded "
-            + "function",
-        )
-        # --- End of dpex passes added to the untyped pipeline               --#
-
-        # inline closures early in case they are using nonlocal's
-        # see issue #6585.
-        pm.add_pass(
-            InlineClosureLikes, "inline calls to locally defined closures"
-        )
-
-        # pre typing
-        if not state.flags.no_rewrites:
-            pm.add_pass(RewriteSemanticConstants, "rewrite semantic constants")
-            pm.add_pass(DeadBranchPrune, "dead branch pruning")
-            pm.add_pass(GenericRewrites, "nopython rewrites")
-
-        # convert any remaining closures into functions
-        pm.add_pass(
-            MakeFunctionToJitFunction,
-            "convert make_function into JIT functions",
-        )
-        # inline functions that have been determined as inlinable and rerun
-        # branch pruning, this needs to be run after closures are inlined as
-        # the IR repr of a closure masks call sites if an inlinable is called
-        # inside a closure
-        pm.add_pass(InlineInlinables, "inline inlinable functions")
-        if not state.flags.no_rewrites:
-            pm.add_pass(DeadBranchPrune, "dead branch pruning")
-
-        pm.add_pass(FindLiterallyCalls, "find literally calls")
-        pm.add_pass(LiteralUnroll, "handles literal_unroll")
-
-        if state.flags.enable_ssa:
-            pm.add_pass(ReconstructSSA, "ssa")
-
-        pm.add_pass(LiteralPropagationSubPipelinePass, "Literal propagation")
-
-        pm.finalize()
-        return pm
-
-    @staticmethod
-    def define_typed_pipeline(state, name="dpex_offload_typed"):
-        """Returns the typed part of the nopython pipeline"""
-        pm = PassManager(name)
-        # typing
-        pm.add_pass(NopythonTypeInference, "nopython frontend")
-        # Annotate only once legalized
-        pm.add_pass(AnnotateTypes, "annotate types")
-        pm.add_pass(
-            RewriteNdarrayFunctionsPass,
-            "Rewrite numpy.ndarray functions to dpnp.ndarray functions",
-        )
-
-        # strip phis
-        pm.add_pass(PreLowerStripPhis, "remove phis nodes")
-
-        # optimization
-        pm.add_pass(InlineOverloads, "inline overloaded functions")
-        pm.add_pass(PreParforPass, "Preprocessing for parfors")
-        if not state.flags.no_rewrites:
-            pm.add_pass(NopythonRewrites, "nopython rewrites")
-        pm.add_pass(ParforPass, "convert to parfors")
-
-        pm.finalize()
-        return pm
-
-    @staticmethod
-    def define_nopython_lowering_pipeline(state, name="dpex_offload_lowering"):
-        """Returns an nopython mode pipeline based PassManager"""
-        pm = PassManager(name)
-
-        # legalize
-        pm.add_pass(
-            NoPythonSupportedFeatureValidation,
-            "ensure features that are in use are in a valid form",
-        )
-        pm.add_pass(IRLegalization, "ensure IR is legal prior to lowering")
-
-        # lower
-        pm.add_pass(DpexLowering, "Custom Lowerer with auto-offload support")
-        pm.add_pass(NoPythonBackend, "nopython mode backend")
-        pm.add_pass(DumpParforDiagnostics, "dump parfor diagnostics")
-
-        pm.finalize()
-        return pm
-
-    @staticmethod
-    def define_nopython_pipeline(state, name="dpex_offload_nopython"):
-        """Returns an nopython mode pipeline based PassManager"""
-        # compose pipeline from untyped, typed and lowering parts
-        dpb = _OffloadPassBuilder
-        pm = PassManager(name)
-        untyped_passes = dpb.define_untyped_pipeline(state)
-        pm.passes.extend(untyped_passes.passes)
-
-        typed_passes = dpb.define_typed_pipeline(state)
-        pm.passes.extend(typed_passes.passes)
-
-        lowering_passes = dpb.define_nopython_lowering_pipeline(state)
-        pm.passes.extend(lowering_passes.passes)
-
-        pm.finalize()
-        return pm
-
-
-class OffloadCompiler(CompilerBase):
-    """Dpex's offload compiler pipeline that can offload parfor nodes and
-    other NumPy-based library calls to a SYCL device.
-
-    .. note:: Deprecated in 0.20
-          The offload compiler is deprecated and will be removed in a future
-          release.
-    """
-
-    def define_pipelines(self):
-        pms = []
-        self.state.parfor_diagnostics = ExtendedParforDiagnostics()
-        self.state.metadata[
-            "parfor_diagnostics"
-        ] = self.state.parfor_diagnostics
-        if not self.state.flags.force_pyobject:
-            pms.append(_OffloadPassBuilder.define_nopython_pipeline(self.state))
-        if self.state.status.can_fallback or self.state.flags.force_pyobject:
-            raise UnsupportedCompilationModeError()
-        return pms
diff --git a/numba_dpex/dpnp_iface/__init__.py b/numba_dpex/dpnp_iface/__init__.py
index 2517018742..d4e46283c8 100644
--- a/numba_dpex/dpnp_iface/__init__.py
+++ b/numba_dpex/dpnp_iface/__init__.py
@@ -3,28 +3,3 @@
 # SPDX-License-Identifier: Apache-2.0
 
 from . import arrayobj
-
-
-def ensure_dpnp(name):
-    try:
-        from . import dpnp_fptr_interface as dpnp_iface
-    except ImportError:
-        raise ImportError("dpnp is needed to call np.%s" % name)
-
-
-def _init_dpnp():
-    try:
-        import os
-
-        import dpnp
-
-        if hasattr(os, "add_dll_directory"):
-            os.add_dll_directory(os.path.dirname(dpnp.__file__))
-    except ImportError:
-        pass
-
-
-_init_dpnp()
-
-
-DEBUG = None
diff --git a/numba_dpex/dpnp_iface/dpnp_array_ops_impl.py b/numba_dpex/dpnp_iface/dpnp_array_ops_impl.py
deleted file mode 100644
index 7be984eed9..0000000000
--- a/numba_dpex/dpnp_iface/dpnp_array_ops_impl.py
+++ /dev/null
@@ -1,344 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import numpy as np
-from numba import types
-from numba.core.extending import overload, register_jitable
-from numba.core.typing import signature
-
-import numba_dpex
-import numba_dpex.dpctl_iface as dpctl_functions
-import numba_dpex.dpnp_iface as dpnp_lowering
-import numba_dpex.dpnp_iface.dpnp_stubs_impl as dpnp_ext
-
-from . import stubs
-
-
-@register_jitable
-def common_impl(a, out, dpnp_func, print_debug):
-    if a.size == 0:
-        raise ValueError("Passed Empty array")
-
-    sycl_queue = dpctl_functions.get_current_queue()
-    a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    out_usm = dpctl_functions.malloc_shared(a.itemsize, sycl_queue)
-
-    dpnp_func(a_usm, out_usm, a.size)
-
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    dpctl_functions.free_with_queue(a_usm, sycl_queue)
-    dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-    dpnp_ext._dummy_liveness_func([a.size, out.size])
-
-    if print_debug:
-        print("dpnp implementation")
-
-
-@overload(stubs.dpnp.cumsum)
-def dpnp_cumsum_impl(a):
-    name = "cumsum"
-    dpnp_lowering.ensure_dpnp(name)
-
-    res_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.5.1/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp#L135
-    Function declaration:
-    void dpnp_cumsum_c(void* array1_in, void* result1, size_t size)
-    """
-    sig = signature(res_type, types.voidptr, types.voidptr, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        out = np.arange(0, a.size, 1, a.dtype)
-        common_impl(a, out, dpnp_func, PRINT_DEBUG)
-
-        return out
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.cumprod)
-def dpnp_cumprod_impl(a):
-    name = "cumprod"
-    dpnp_lowering.ensure_dpnp(name)
-
-    res_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.5.1/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp#L110
-    Function declaration:
-    void dpnp_cumprod_c(void* array1_in, void* result1, size_t size)
-    """
-    sig = signature(res_type, types.voidptr, types.voidptr, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-    if a.dtype == types.Integer:
-        ret_dtype = np.int64
-    else:
-        ret_dtype = a.dtype
-
-    def dpnp_impl(a):
-        out = np.arange(0, a.size, 1, ret_dtype)
-        common_impl(a, out, dpnp_func, PRINT_DEBUG)
-
-        return out
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.copy)
-def dpnp_copy_impl(a):
-    name = "copy"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.10.0dev0/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp#L320-L330
-    Function declaration:
-    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);
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,  # void* result_out
-        types.intp,  # const size_t result_size,
-        types.intp,  # const size_t result_ndim,
-        types.voidptr,  # const shape_elem_type* result_shape,
-        types.voidptr,  # const shape_elem_type* result_strides,
-        types.voidptr,  # const void* input1_in,
-        types.intp,  # const size_t input1_size,
-        types.intp,  # const size_t input1_ndim,
-        types.voidptr,  # const shape_elem_type* input1_shape,
-        types.voidptr,  # const shape_elem_type* input1_strides,
-        types.voidptr,  # const size_t* where);
-    )
-
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-
-    res_dtype = a.dtype
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        if a.size == 0:
-            raise ValueError("Passed Empty array")
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out = np.arange(0, a.size, 1, res_dtype)
-        out_usm = dpctl_functions.malloc_shared(
-            out.size * out.itemsize, sycl_queue
-        )
-
-        strides = np.array(1)
-
-        result_out = out_usm
-        result_size = out.size
-        result_ndim = out.ndim
-        result_shape = out.shapeptr
-        result_strides = strides.ctypes
-
-        input1_in = a_usm
-        input1_size = a.size
-        input1_ndim = a.ndim
-        input1_shape = a.shapeptr
-        input1_strides = strides.ctypes
-
-        where = 0
-
-        dpnp_func(
-            result_out,
-            result_size,
-            result_ndim,
-            result_shape,
-            result_strides,
-            input1_in,
-            input1_size,
-            input1_ndim,
-            input1_shape,
-            input1_strides,
-            where,
-        )
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a.size, out.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        return out
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.sort)
-def dpnp_sort_impl(a):
-    name = "sort"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.5.0/dpnp/backend/kernels/dpnp_krnl_sorting.cpp#L90
-
-    Function declaration:
-    void dpnp_sort_c(void* array1_in, void* result1, size_t size)
-
-    """
-    sig = signature(ret_type, types.voidptr, types.voidptr, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-
-    res_dtype = a.dtype
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        if a.size == 0:
-            raise ValueError("Passed Empty array")
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out = np.arange(0, a.size, 1, res_dtype)
-        out_usm = dpctl_functions.malloc_shared(
-            out.size * out.itemsize, sycl_queue
-        )
-
-        dpnp_func(a_usm, out_usm, a.size)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a.size, out.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        return out
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.take)
-def dpnp_take_impl(a, ind):
-    name = "take"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/9b14f0ca76a9e0c309bb97b4d5caa0870eecd6bb/dpnp/backend/kernels/dpnp_krnl_indexing.cpp#L925
-    Function declaration:
-    void dpnp_take_c(void* array1_in, const size_t array1_size, void* indices1, void* result1, size_t size)
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,
-        types.intp,
-        types.voidptr,
-        types.voidptr,
-        types.intp,
-    )
-    dpnp_func = dpnp_ext.dpnp_func(
-        "dpnp_" + name, [a.dtype.name, ind.dtype.name], sig
-    )
-
-    res_dtype = a.dtype
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a, ind):
-        if a.size == 0:
-            raise ValueError("Passed Empty array")
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        ind_usm = dpctl_functions.malloc_shared(
-            ind.size * ind.itemsize, sycl_queue
-        )
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, ind_usm, ind.ctypes, ind.size * ind.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out = np.arange(0, ind.size, 1, res_dtype).reshape(ind.shape)
-        out_usm = dpctl_functions.malloc_shared(
-            out.size * out.itemsize, sycl_queue
-        )
-
-        dpnp_func(a_usm, a.size * a.itemsize, ind_usm, out_usm, ind.size)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(ind_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a.size, ind.size, out.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        return out
-
-    return dpnp_impl
diff --git a/numba_dpex/dpnp_iface/dpnp_fptr_interface.pyx b/numba_dpex/dpnp_iface/dpnp_fptr_interface.pyx
deleted file mode 100644
index 541a8da5ff..0000000000
--- a/numba_dpex/dpnp_iface/dpnp_fptr_interface.pyx
+++ /dev/null
@@ -1,267 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2022 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-# distutils: language = c++
-# cython: language_level=3
-
-import ctypes
-
-
-cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName":  # need this namespace for Enum import
-    cdef enum DPNPFuncName "DPNPFuncName":
-        DPNP_FN_ABSOLUTE
-        DPNP_FN_ADD
-        DPNP_FN_ALL
-        DPNP_FN_ARANGE
-        DPNP_FN_ARCCOS
-        DPNP_FN_ARCCOSH
-        DPNP_FN_ARCSIN
-        DPNP_FN_ARCSINH
-        DPNP_FN_ARCTAN
-        DPNP_FN_ARCTAN2
-        DPNP_FN_ARCTANH
-        DPNP_FN_ARGMAX
-        DPNP_FN_ARGMIN
-        DPNP_FN_ARGSORT
-        DPNP_FN_BITWISE_AND
-        DPNP_FN_BITWISE_OR
-        DPNP_FN_BITWISE_XOR
-        DPNP_FN_CBRT
-        DPNP_FN_CEIL
-        DPNP_FN_CHOLESKY
-        DPNP_FN_COPY
-        DPNP_FN_COPYSIGN
-        DPNP_FN_CORRELATE
-        DPNP_FN_COS
-        DPNP_FN_COSH
-        DPNP_FN_COV
-        DPNP_FN_CUMPROD
-        DPNP_FN_CUMSUM
-        DPNP_FN_DEGREES
-        DPNP_FN_DET
-        DPNP_FN_DIAGONAL
-        DPNP_FN_DIVIDE
-        DPNP_FN_DOT
-        DPNP_FN_EIG
-        DPNP_FN_EIGVALS
-        DPNP_FN_EXP
-        DPNP_FN_EXP2
-        DPNP_FN_EXPM1
-        DPNP_FN_FABS
-        DPNP_FN_FFT_FFT
-        DPNP_FN_FLOOR
-        DPNP_FN_FLOOR_DIVIDE
-        DPNP_FN_FMOD
-        DPNP_FN_HYPOT
-        DPNP_FN_INITVAL
-        DPNP_FN_INVERT
-        DPNP_FN_LEFT_SHIFT
-        DPNP_FN_LOG
-        DPNP_FN_LOG10
-        DPNP_FN_LOG1P
-        DPNP_FN_LOG2
-        DPNP_FN_MATMUL
-        DPNP_FN_MATRIX_RANK
-        DPNP_FN_MAX
-        DPNP_FN_MAXIMUM
-        DPNP_FN_MEAN
-        DPNP_FN_MEDIAN
-        DPNP_FN_MIN
-        DPNP_FN_MINIMUM
-        DPNP_FN_MODF
-        DPNP_FN_MULTIPLY
-        DPNP_FN_PARTITION
-        DPNP_FN_POWER
-        DPNP_FN_PROD
-        DPNP_FN_RADIANS
-        DPNP_FN_RECIP
-        DPNP_FN_REMAINDER
-        DPNP_FN_REPEAT
-        DPNP_FN_RIGHT_SHIFT
-        DPNP_FN_RNG_BETA
-        DPNP_FN_RNG_BINOMIAL
-        DPNP_FN_RNG_CHISQUARE
-        DPNP_FN_RNG_EXPONENTIAL
-        DPNP_FN_RNG_GAMMA
-        DPNP_FN_RNG_GAUSSIAN
-        DPNP_FN_RNG_GEOMETRIC
-        DPNP_FN_RNG_GUMBEL
-        DPNP_FN_RNG_HYPERGEOMETRIC
-        DPNP_FN_RNG_LAPLACE
-        DPNP_FN_RNG_LOGNORMAL
-        DPNP_FN_RNG_MULTINOMIAL
-        DPNP_FN_RNG_MULTIVARIATE_NORMAL
-        DPNP_FN_RNG_NEGATIVE_BINOMIAL
-        DPNP_FN_RNG_NORMAL
-        DPNP_FN_RNG_POISSON
-        DPNP_FN_RNG_RAYLEIGH
-        DPNP_FN_RNG_STANDARD_CAUCHY
-        DPNP_FN_RNG_STANDARD_EXPONENTIAL
-        DPNP_FN_RNG_STANDARD_GAMMA
-        DPNP_FN_RNG_STANDARD_NORMAL
-        DPNP_FN_RNG_UNIFORM
-        DPNP_FN_RNG_WEIBULL
-        DPNP_FN_SIGN
-        DPNP_FN_SIN
-        DPNP_FN_SINH
-        DPNP_FN_SORT
-        DPNP_FN_SQRT
-        DPNP_FN_SQUARE
-        DPNP_FN_STD
-        DPNP_FN_SUBTRACT
-        DPNP_FN_SUM
-        DPNP_FN_TAKE
-        DPNP_FN_TAN
-        DPNP_FN_TANH
-        DPNP_FN_TRANSPOSE
-        DPNP_FN_TRUNC
-        DPNP_FN_VAR
-
-
-cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncType":  # need this namespace for Enum import
-    cdef enum DPNPFuncType "DPNPFuncType":
-        DPNP_FT_NONE
-        DPNP_FT_INT
-        DPNP_FT_LONG
-        DPNP_FT_FLOAT
-        DPNP_FT_DOUBLE
-        DPNP_FT_BOOL
-
-cdef extern from "dpnp_iface_fptr.hpp":
-    struct DPNPFuncData:
-        DPNPFuncType return_type
-        void * ptr
-
-    DPNPFuncData get_dpnp_function_ptr(DPNPFuncName name, DPNPFuncType first_type, DPNPFuncType second_type) except +
-
-
-
-cdef DPNPFuncName get_DPNPFuncName_from_str(name):
-    if name == "dpnp_all":
-        return DPNPFuncName.DPNP_FN_ALL
-    if name == "dpnp_argmax":
-        return DPNPFuncName.DPNP_FN_ARGMAX
-    if name == "dpnp_argmin":
-        return DPNPFuncName.DPNP_FN_ARGMIN
-    if name == "dpnp_argsort":
-        return DPNPFuncName.DPNP_FN_ARGSORT
-    if name == "dpnp_beta":
-        return DPNPFuncName.DPNP_FN_RNG_BETA
-    if name == "dpnp_binomial":
-        return DPNPFuncName.DPNP_FN_RNG_BINOMIAL
-    if name == "dpnp_chisquare":
-        return DPNPFuncName.DPNP_FN_RNG_CHISQUARE
-    if name == "dpnp_cholesky":
-        return DPNPFuncName.DPNP_FN_CHOLESKY
-    if name == "dpnp_copy":
-        return DPNPFuncName.DPNP_FN_COPY
-    if name == "dpnp_cov":
-        return DPNPFuncName.DPNP_FN_COV
-    if name == "dpnp_cumprod":
-        return DPNPFuncName.DPNP_FN_CUMPROD
-    if name == "dpnp_cumsum":
-        return DPNPFuncName.DPNP_FN_CUMSUM
-    if name == "dpnp_det":
-        return DPNPFuncName.DPNP_FN_DET
-    if name == "dpnp_diagonal":
-        return DPNPFuncName.DPNP_FN_DIAGONAL
-    if name == "dpnp_dot":
-        return DPNPFuncName.DPNP_FN_DOT
-    if name == "dpnp_eig":
-        return DPNPFuncName.DPNP_FN_EIG
-    if name == "dpnp_exponential":
-        return DPNPFuncName.DPNP_FN_RNG_EXPONENTIAL
-    if name == "dpnp_gamma":
-        return DPNPFuncName.DPNP_FN_RNG_GAMMA
-    if name == "dpnp_geometric":
-        return DPNPFuncName.DPNP_FN_RNG_GEOMETRIC
-    if name == "dpnp_gumbel":
-        return DPNPFuncName.DPNP_FN_RNG_GUMBEL
-    if name == "dpnp_hypergeometric":
-        return DPNPFuncName.DPNP_FN_RNG_HYPERGEOMETRIC
-    if name == "dpnp_laplace":
-        return DPNPFuncName.DPNP_FN_RNG_LAPLACE
-    if name == "dpnp_lognormal":
-        return DPNPFuncName.DPNP_FN_RNG_LOGNORMAL
-    if name == "dpnp_matmul":
-        return DPNPFuncName.DPNP_FN_MATMUL
-    if name == "dpnp_matrix_rank":
-        return DPNPFuncName.DPNP_FN_MATRIX_RANK
-    if name == "dpnp_max":
-        return DPNPFuncName.DPNP_FN_MAX
-    if name == "dpnp_mean":
-        return DPNPFuncName.DPNP_FN_MEAN
-    if name == "dpnp_median":
-        return DPNPFuncName.DPNP_FN_MEDIAN
-    if name == "dpnp_min":
-        return DPNPFuncName.DPNP_FN_MIN
-    if name == "dpnp_multinomial":
-        return DPNPFuncName.DPNP_FN_RNG_MULTINOMIAL
-    if name == "dpnp_multivariate_normal":
-        return DPNPFuncName.DPNP_FN_RNG_MULTIVARIATE_NORMAL
-    if name == "dpnp_negative_binomial":
-        return DPNPFuncName.DPNP_FN_RNG_NEGATIVE_BINOMIAL
-    if name == "dpnp_normal":
-        return DPNPFuncName.DPNP_FN_RNG_NORMAL
-    if name == "dpnp_partition":
-        return DPNPFuncName.DPNP_FN_PARTITION
-    if name == "dpnp_poisson":
-        return DPNPFuncName.DPNP_FN_RNG_POISSON
-    if name == "dpnp_prod":
-        return DPNPFuncName.DPNP_FN_PROD
-    if name == "dpnp_random_sample":
-        return DPNPFuncName.DPNP_FN_RNG_UNIFORM
-    if name == "dpnp_rayleigh":
-        return DPNPFuncName.DPNP_FN_RNG_RAYLEIGH
-    if name == "dpnp_repeat":
-        return DPNPFuncName.DPNP_FN_REPEAT
-    if name == "dpnp_sort":
-        return DPNPFuncName.DPNP_FN_SORT
-    if name == "dpnp_standard_cauchy":
-        return DPNPFuncName.DPNP_FN_RNG_STANDARD_CAUCHY
-    if name == "dpnp_standard_exponential":
-        return DPNPFuncName.DPNP_FN_RNG_STANDARD_EXPONENTIAL
-    if name == "dpnp_standard_gamma":
-        return DPNPFuncName.DPNP_FN_RNG_STANDARD_GAMMA
-    if name == "dpnp_sum":
-        return DPNPFuncName.DPNP_FN_SUM
-    if name == "dpnp_take":
-        return DPNPFuncName.DPNP_FN_TAKE
-    if name == "dpnp_uniform":
-        return DPNPFuncName.DPNP_FN_RNG_UNIFORM
-    if name == "dpnp_vdot":
-        return DPNPFuncName.DPNP_FN_DOT
-    if name == "dpnp_weibull":
-        return DPNPFuncName.DPNP_FN_RNG_WEIBULL
-
-    raise ValueError("Unknown dpnp function requested: " + name.split("_")[1])
-
-
-cdef DPNPFuncType get_DPNPFuncType_from_str(name):
-    if name == "float32":
-        return DPNPFuncType.DPNP_FT_FLOAT
-    elif name == "int32" or name == "uint32":
-        return DPNPFuncType.DPNP_FT_INT
-    elif name == "float64":
-        return DPNPFuncType.DPNP_FT_DOUBLE
-    elif name == "int64" or name == "uint64":
-        return DPNPFuncType.DPNP_FT_LONG
-    elif name == "bool":
-        return DPNPFuncType.DPNP_FT_BOOL
-    else:
-        return DPNPFuncType.DPNP_FT_NONE
-
-from libc.stdint cimport uintptr_t
-from libc.stdio cimport printf
-
-
-cpdef get_dpnp_fn_ptr(name, types):
-    cdef DPNPFuncName func_name = get_DPNPFuncName_from_str(name)
-    cdef DPNPFuncType first_type = get_DPNPFuncType_from_str(types[0])
-    cdef DPNPFuncType second_type = get_DPNPFuncType_from_str(types[1])
-
-    cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(func_name, first_type, second_type)
-    cdef uintptr_t fn_ptr = <uintptr_t>kernel_data.ptr
-
-    return <object>fn_ptr
diff --git a/numba_dpex/dpnp_iface/dpnp_indexing.py b/numba_dpex/dpnp_iface/dpnp_indexing.py
deleted file mode 100644
index fb85cd528a..0000000000
--- a/numba_dpex/dpnp_iface/dpnp_indexing.py
+++ /dev/null
@@ -1,115 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import numpy as np
-from numba import types
-from numba.core.extending import overload, register_jitable
-from numba.core.typing import signature
-
-import numba_dpex.dpctl_iface as dpctl_functions
-import numba_dpex.dpnp_iface as dpnp_lowering
-import numba_dpex.dpnp_iface.dpnp_stubs_impl as dpnp_ext
-
-from . import stubs
-
-
-@overload(stubs.dpnp.diagonal)
-def dpnp_diagonal_impl(a, offset=0):
-    name = "diagonal"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/e389248c709531b181be8bf33b1a270fca812a92/dpnp/backend/kernels/dpnp_krnl_indexing.cpp#L39
-
-    Function declaration:
-    void dpnp_diagonal_c(
-        void* array1_in, const size_t input1_size, void* result1, const size_t offset, size_t* shape, size_t* res_shape, const size_t res_ndim)
-
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,
-        types.intp,
-        types.voidptr,
-        types.intp,
-        types.voidptr,
-        types.voidptr,
-        types.intp,
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    function_text = f"""\
-def tuplizer(a):
-    return ({", ".join(f"a[{i}]" for i in range(a.ndim - 1))})
-"""
-    locals = {}
-    exec(function_text, globals(), locals)
-    tuplizer = register_jitable(locals["tuplizer"])
-
-    def dpnp_impl(a, offset=0):
-        if a.size == 0:
-            raise ValueError("Passed Empty array")
-
-        n = min(a.shape[0], a.shape[1])
-        res_shape = np.zeros(a.ndim - 1, dtype=np.int64)
-
-        if a.ndim > 2:
-            for i in range(a.ndim - 2):
-                res_shape[i] = a.shape[i + 2]
-
-        if (n + offset) > a.shape[1]:
-            res_shape[-1] = a.shape[1] - offset
-        elif (n + offset) > a.shape[0]:
-            res_shape[-1] = a.shape[0]
-        else:
-            res_shape[-1] = n + offset
-
-        shape = tuplizer(res_shape)
-
-        out = np.empty(shape, dtype=a.dtype)
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out_usm = dpctl_functions.malloc_shared(
-            out.size * out.itemsize, sycl_queue
-        )
-
-        dpnp_func(
-            a_usm,
-            a.size * a.itemsize,
-            out_usm,
-            offset,
-            a.shapeptr,
-            out.shapeptr,
-            out.ndim,
-        )
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a.size, out.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-
-        return out
-
-    return dpnp_impl
diff --git a/numba_dpex/dpnp_iface/dpnp_linalgimpl.py b/numba_dpex/dpnp_iface/dpnp_linalgimpl.py
deleted file mode 100644
index 982d360757..0000000000
--- a/numba_dpex/dpnp_iface/dpnp_linalgimpl.py
+++ /dev/null
@@ -1,734 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import numpy as np
-from numba import types
-from numba.core.extending import overload, register_jitable
-from numba.core.typing import signature
-
-import numba_dpex
-import numba_dpex.dpctl_iface as dpctl_functions
-import numba_dpex.dpnp_iface as dpnp_lowering
-import numba_dpex.dpnp_iface.dpnp_stubs_impl as dpnp_ext
-
-from . import stubs
-
-
-@overload(stubs.dpnp.eig)
-def dpnp_eig_impl(a):
-    name = "eig"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels.cpp#L180
-
-    Function declaration:
-    void dpnp_eig_c(const void* array_in, void* result1, void* result2, size_t size)
-
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.voidptr, types.voidptr, types.intp
-    )
-    dpnp_eig = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-
-    res_dtype = np.float64
-    if a.dtype == types.float32:
-        res_dtype = np.float32
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        n = a.shape[-1]
-        if a.shape[-2] != n:
-            msg = "Last 2 dimensions of the array must be square."
-            raise ValueError(msg)
-
-        dpnp_ext._check_finite_matrix(a)
-
-        wr = np.empty(n, dtype=res_dtype)
-        vr = np.empty((n, n), dtype=res_dtype)
-
-        if n == 0:
-            return (wr, vr)
-
-        sycl_queue = dpctl_functions.get_current_queue()
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        wr_usm = dpctl_functions.malloc_shared(
-            wr.size * wr.itemsize, sycl_queue
-        )
-        vr_usm = dpctl_functions.malloc_shared(
-            vr.size * vr.itemsize, sycl_queue
-        )
-
-        dpnp_eig(a_usm, wr_usm, vr_usm, n)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, wr.ctypes, wr_usm, wr.size * wr.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, vr.ctypes, vr_usm, vr.size * vr.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(wr_usm, sycl_queue)
-        dpctl_functions.free_with_queue(vr_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([wr.size, vr.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        return (wr, vr)
-
-    return dpnp_impl
-
-
-@register_jitable
-def common_matmul_impl(dpnp_func, a, b, out, m, n, k, print_debug):
-    sycl_queue = dpctl_functions.get_current_queue()
-
-    a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-    dpctl_functions.queue_memcpy(
-        sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-    )
-
-    b_usm = dpctl_functions.malloc_shared(b.size * b.itemsize, sycl_queue)
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, b_usm, b.ctypes, b.size * b.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    out_usm = dpctl_functions.malloc_shared(out.size * out.itemsize, sycl_queue)
-
-    a_m = np.array((m, k))
-    b_m = np.array((k, n))
-    out_m = np.array((m, n))
-
-    strides = np.array(1)
-
-    result_out = out_usm
-    result_size = out.size
-    result_ndim = 2
-    result_shape = out_m.ctypes
-    result_strides = strides.ctypes
-
-    input1_in = a_usm
-    input1_size = a.size
-    input1_ndim = 2
-    input1_shape = a_m.ctypes
-    input1_strides = strides.ctypes
-
-    input2_in = b_usm
-    input2_size = b.size
-    input2_ndim = 2
-    input2_shape = b_m.ctypes
-    input2_strides = strides.ctypes
-
-    dpnp_func(
-        result_out,
-        result_size,
-        result_ndim,
-        result_shape,
-        result_strides,
-        input1_in,
-        input1_size,
-        input1_ndim,
-        input1_shape,
-        input1_strides,
-        input2_in,
-        input2_size,
-        input2_ndim,
-        input2_shape,
-        input2_strides,
-    )
-
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    dpctl_functions.free_with_queue(a_usm, sycl_queue)
-    dpctl_functions.free_with_queue(b_usm, sycl_queue)
-    dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-    dpnp_ext._dummy_liveness_func([a.size, b.size, out.size])
-    dpnp_ext._dummy_liveness_func([a_m.size, b_m.size, out_m.size])
-
-    if print_debug:
-        print("dpnp implementation")
-
-
-@register_jitable
-def common_dot_impl(dpnp_func, a, b, out, m, print_debug):
-    sycl_queue = dpctl_functions.get_current_queue()
-    a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    b_usm = dpctl_functions.malloc_shared(b.size * b.itemsize, sycl_queue)
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, b_usm, b.ctypes, b.size * b.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    out_usm = dpctl_functions.malloc_shared(out.size * out.itemsize, sycl_queue)
-
-    strides = np.array(1)
-
-    result_out = out_usm
-    result_size = out.size
-    result_ndim = out.ndim
-    result_shape = out.shapeptr
-    result_strides = strides.ctypes
-
-    input1_in = a_usm
-    input1_size = a.size
-    input1_ndim = a.ndim
-    input1_shape = a.shapeptr
-    input1_strides = strides.ctypes
-
-    input2_in = b_usm
-    input2_size = b.size
-    input2_ndim = b.ndim
-    input2_shape = b.shapeptr
-    input2_strides = strides.ctypes
-
-    dpnp_func(
-        result_out,
-        result_size,
-        result_ndim,
-        result_shape,
-        result_strides,
-        input1_in,
-        input1_size,
-        input1_ndim,
-        input1_shape,
-        input1_strides,
-        input2_in,
-        input2_size,
-        input2_ndim,
-        input2_shape,
-        input2_strides,
-    )
-
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    dpctl_functions.free_with_queue(a_usm, sycl_queue)
-    dpctl_functions.free_with_queue(b_usm, sycl_queue)
-    dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-    dpnp_ext._dummy_liveness_func([a.size, b.size, out.size])
-
-    if print_debug:
-        print("dpnp implementation")
-
-
-def get_res_dtype(a, b):
-    res_dtype = np.float64
-    if a.dtype == types.int32 and b.dtype == types.int32:
-        res_dtype = np.int32
-    elif a.dtype == types.int32 and b.dtype == types.int64:
-        res_dtype = np.int64
-    elif a.dtype == types.int32 and b.dtype == types.float32:
-        res_dtype = np.float64
-    elif a.dtype == types.int32 and b.dtype == types.float64:
-        res_dtype = np.float64
-    elif a.dtype == types.int64 and b.dtype == types.int32:
-        res_dtype = np.int64
-    elif a.dtype == types.int64 and b.dtype == types.int64:
-        res_dtype = np.int64
-    elif a.dtype == types.int64 and b.dtype == types.float32:
-        res_dtype = np.float64
-    elif a.dtype == types.int64 and b.dtype == types.float64:
-        res_dtype = np.float64
-    elif a.dtype == types.float32 and b.dtype == types.int32:
-        res_dtype = np.float64
-    elif a.dtype == types.float32 and b.dtype == types.int64:
-        res_dtype = np.float64
-    elif a.dtype == types.float32 and b.dtype == types.float32:
-        res_dtype = np.float32
-    elif a.dtype == types.float32 and b.dtype == types.float64:
-        res_dtype = np.float64
-    elif a.dtype == types.float64 and b.dtype == types.int32:
-        res_dtype = np.float64
-    elif a.dtype == types.float64 and b.dtype == types.int64:
-        res_dtype = np.float64
-    elif a.dtype == types.float64 and b.dtype == types.float32:
-        res_dtype = np.float32
-    elif a.dtype == types.float64 and b.dtype == types.float64:
-        res_dtype = np.float64
-
-    return res_dtype
-
-
-@overload(stubs.dpnp.matmul)
-@overload(stubs.dpnp.dot)
-def dpnp_dot_impl(a, b):
-    dpnp_lowering.ensure_dpnp("dot")
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blame/67a101c90cf253cfe9b9ba80ac397811ce94edee/dpnp/backend/kernels/dpnp_krnl_common.cpp#L322
-
-    Function declaration:
-    void dpnp_matmul_c(void* result_out,
-                    const size_t result_size,
-                    const size_t result_ndim,
-                    const size_t* result_shape,
-                    const size_t* result_strides,
-                    const void* input1_in,
-                    const size_t input1_size,
-                    const size_t input1_ndim,
-                    const size_t* input1_shape,
-                    const size_t* input1_strides,
-                    const void* input2_in,
-                    const size_t input2_size,
-                    const size_t input2_ndim,
-                    const size_t* input2_shape,
-                    const size_t* input2_strides)
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,
-        types.intp,
-        types.intp,
-        types.voidptr,
-        types.voidptr,
-        types.voidptr,
-        types.intp,
-        types.intp,
-        types.voidptr,
-        types.voidptr,
-        types.voidptr,
-        types.intp,
-        types.intp,
-        types.voidptr,
-        types.voidptr,
-    )
-
-    res_dtype = get_res_dtype(a, b)
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    ndims = [a.ndim, b.ndim]
-    if ndims == [2, 2]:
-        dpnp_func = dpnp_ext.dpnp_func(
-            "dpnp_matmul", [a.dtype.name, "NONE"], sig
-        )
-
-        def dot_2_mm(a, b):
-            m, k = a.shape
-            _k, n = b.shape
-
-            if _k != k:
-                raise ValueError("Incompatible array sizes for np.dot(a, b)")
-
-            out = np.empty((m, n), dtype=res_dtype)
-            common_matmul_impl(dpnp_func, a, b, out, m, n, k, PRINT_DEBUG)
-
-            return out
-
-        return dot_2_mm
-    elif ndims == [2, 1]:
-        dpnp_func = dpnp_ext.dpnp_func(
-            "dpnp_matmul", [a.dtype.name, "NONE"], sig
-        )
-
-        def dot_2_mv(a, b):
-            m, k = a.shape
-            (_n,) = b.shape
-            n = 1
-
-            if _n != k:
-                raise ValueError("Incompatible array sizes for np.dot(a, b)")
-
-            out = np.empty((m,), dtype=res_dtype)
-            common_matmul_impl(dpnp_func, a, b, out, m, n, k, PRINT_DEBUG)
-
-            return out
-
-        return dot_2_mv
-    elif ndims == [1, 2]:
-        dpnp_func = dpnp_ext.dpnp_func(
-            "dpnp_matmul", [a.dtype.name, "NONE"], sig
-        )
-
-        def dot_2_vm(a, b):
-            (m,) = a.shape
-            k, n = b.shape
-
-            if m != k:
-                raise ValueError("Incompatible array sizes for np.dot(a, b)")
-
-            out = np.empty((n,), dtype=res_dtype)
-            common_matmul_impl(dpnp_func, a, b, out, m, n, k, PRINT_DEBUG)
-
-            return out
-
-        return dot_2_vm
-    elif ndims == [1, 1]:
-        """
-        dpnp source:
-        https://github.com/IntelPython/dpnp/blob/67a101c90cf253cfe9b9ba80ac397811ce94edee/dpnp/backend/kernels/dpnp_krnl_common.cpp#L79
-
-        Function declaration:
-        void dpnp_dot_c(void* result_out,
-                        const size_t result_size,
-                        const size_t result_ndim,
-                        const size_t* result_shape,
-                        const size_t* result_strides,
-                        const void* input1_in,
-                        const size_t input1_size,
-                        const size_t input1_ndim,
-                        const size_t* input1_shape,
-                        const size_t* input1_strides,
-                        const void* input2_in,
-                        const size_t input2_size,
-                        const size_t input2_ndim,
-                        const size_t* input2_shape,
-                        const size_t* input2_strides)
-        """
-        sig = signature(
-            ret_type,
-            types.voidptr,
-            types.intp,
-            types.intp,
-            types.voidptr,
-            types.voidptr,
-            types.voidptr,
-            types.intp,
-            types.intp,
-            types.voidptr,
-            types.voidptr,
-            types.voidptr,
-            types.intp,
-            types.intp,
-            types.voidptr,
-            types.voidptr,
-        )
-        dpnp_func = dpnp_ext.dpnp_func("dpnp_dot", [a.dtype.name, "NONE"], sig)
-
-        def dot_2_vv(a, b):
-            (m,) = a.shape
-            (n,) = b.shape
-
-            if m != n:
-                raise ValueError("Incompatible array sizes for np.dot(a, b)")
-
-            out = np.empty(1, dtype=res_dtype)
-            common_dot_impl(dpnp_func, a, b, out, m, PRINT_DEBUG)
-
-            return out[0]
-
-        return dot_2_vv
-    else:
-        assert 0
-
-
-@overload(stubs.dpnp.multi_dot)
-def dpnp_multi_dot_impl(arrays):
-    dpnp_lowering.ensure_dpnp("multi_dot")
-
-    print_debug = dpnp_lowering.DEBUG
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels.cpp#L118
-
-    Function declaration:
-    void dpnp_dot_c(void* array1_in, void* array2_in, void* result1, size_t size)
-
-    """
-
-    def dpnp_impl(arrays):
-        n = len(arrays)
-        result = arrays[0]
-
-        for idx in range(1, n):
-            result = numba_dpex.dpnp.dot(result, arrays[idx])
-
-        if print_debug:
-            print("dpnp implementation")
-        return result
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.vdot)
-def dpnp_vdot_impl(a, b):
-    dpnp_lowering.ensure_dpnp("vdot")
-
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels.cpp#L118
-
-    Function declaration:
-    void dpnp_dot_c(void* array1_in, void* array2_in, void* result1, size_t size)
-
-    """
-
-    def dpnp_impl(a, b):
-        return numba_dpex.dpnp.dot(np.ravel(a), np.ravel(b))
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.matrix_power)
-def dpnp_matrix_power_impl(a, n):
-    dpnp_lowering.ensure_dpnp("matrix_power")
-
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels.cpp#L42
-
-    Function declaration:
-    void dpnp_matmul_c(void* array1_in, void* array2_in, void* result1, size_t size_m,
-                       size_t size_n, size_t size_k)
-    """
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a, n):
-        if n < 0:
-            raise ValueError(
-                "n < 0 is not supported for np.linalg.matrix_power(a, n)"
-            )
-
-        if n == 0:
-            if PRINT_DEBUG:
-                print("dpnp implementation")
-            return np.identity(a.shape[0], a.dtype)
-
-        result = a
-        for idx in range(0, n - 1):
-            result = numba_dpex.dpnp.matmul(result, a)
-        return result
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.cholesky)
-def dpnp_cholesky_impl(a):
-    name = "cholesky"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_linalg.cpp#L40
-
-    Function declaration:
-    void custom_cholesky_c(void* array1_in, void* result1, size_t* shape)
-
-    """
-    sig = signature(ret_type, types.voidptr, types.voidptr, types.voidptr)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        n = a.shape[-1]
-        if a.shape[-2] != n:
-            raise ValueError("Input array must be square.")
-
-        out = a.copy()
-
-        if n == 0:
-            return out
-
-        sycl_queue = dpctl_functions.get_current_queue()
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out_usm = dpctl_functions.malloc_shared(
-            out.size * out.itemsize, sycl_queue
-        )
-
-        dpnp_func(a_usm, out_usm, a.shapeptr)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([out.size, a.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        return out
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.det)
-def dpnp_det_impl(a):
-    name = "det"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_linalg.cpp#L83
-
-    Function declaration:
-    void custom_det_c(void* array1_in, void* result1, size_t* shape, size_t ndim)
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.voidptr, types.voidptr, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        n = a.shape[-1]
-        if a.shape[-2] != n:
-            raise ValueError("Input array must be square.")
-
-        dpnp_ext._check_finite_matrix(a)
-
-        if a.ndim == 2:
-            out = np.empty((1,), dtype=a.dtype)
-            out[0] = -4
-        else:
-            out = np.empty(a.shape[:-2], dtype=a.dtype)
-
-        sycl_queue = dpctl_functions.get_current_queue()
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out_usm = dpctl_functions.malloc_shared(
-            out.size * out.itemsize, sycl_queue
-        )
-
-        dpnp_func(a_usm, out_usm, a.shapeptr, a.ndim)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([out.size, a.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        if a.ndim == 2:
-            return out[0]
-        else:
-            return out
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.matrix_rank)
-def dpnp_matrix_rank_impl(M, tol=None, hermitian=False):
-    name = "matrix_rank"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_linalg.cpp#L186
-
-    Function declaration:
-    void custom_matrix_rank_c(void* array1_in, void* result1, size_t* shape, size_t ndim)
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.voidptr, types.voidptr, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [M.dtype.name, "NONE"], sig)
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(M, tol=None, hermitian=False):
-        if tol is not None:
-            raise ValueError(
-                "tol is not supported for np.linalg.matrix_rank(M)"
-            )
-        if hermitian:
-            raise ValueError(
-                "hermitian is not supported for np.linalg.matrix_rank(M)"
-            )
-
-        if M.ndim > 2:
-            raise ValueError(
-                "np.linalg.matrix_rank(M) is only supported on 1 or 2-d arrays"
-            )
-
-        out = np.empty(1, dtype=M.dtype)
-
-        sycl_queue = dpctl_functions.get_current_queue()
-        M_usm = dpctl_functions.malloc_shared(M.size * M.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, M_usm, M.ctypes, M.size * M.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out_usm = dpctl_functions.malloc_shared(
-            out.size * out.itemsize, sycl_queue
-        )
-
-        dpnp_func(M_usm, out_usm, M.shapeptr, M.ndim)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(M_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([out.size, M.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        return out[0]
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.eigvals)
-def dpnp_eigvals_impl(a):
-    dpnp_lowering.ensure_dpnp("eigvals")
-
-    def dpnp_impl(a):
-        eigval, eigvec = numba_dpex.dpnp.eig(a)
-        return eigval
-
-    return dpnp_impl
diff --git a/numba_dpex/dpnp_iface/dpnp_logic.py b/numba_dpex/dpnp_iface/dpnp_logic.py
deleted file mode 100644
index db5740110b..0000000000
--- a/numba_dpex/dpnp_iface/dpnp_logic.py
+++ /dev/null
@@ -1,72 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import numpy as np
-from numba import types
-from numba.core.extending import overload
-from numba.core.typing import signature
-
-import numba_dpex.dpctl_iface as dpctl_functions
-import numba_dpex.dpnp_iface as dpnp_lowering
-import numba_dpex.dpnp_iface.dpnp_stubs_impl as dpnp_ext
-
-from . import stubs
-
-
-@overload(stubs.dpnp.all)
-def dpnp_all_impl(a):
-    name = "all"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.6.2/dpnp/backend/kernels/dpnp_krnl_logic.cpp#L36
-    Function declaration:
-    void dpnp_all_c(const void* array1_in, void* result1, const size_t size)
-    """
-    sig = signature(ret_type, types.voidptr, types.voidptr, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        if a.size == 0:
-            return True
-
-        out = np.empty(1, dtype=np.bool_)
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out_usm = dpctl_functions.malloc_shared(
-            out.size * out.itemsize, sycl_queue
-        )
-
-        dpnp_func(a_usm, out_usm, a.size)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a.size, out.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-
-        # TODO: sometimes all() returns ndarray
-        return out[0]
-
-    return dpnp_impl
diff --git a/numba_dpex/dpnp_iface/dpnp_manipulation.py b/numba_dpex/dpnp_iface/dpnp_manipulation.py
deleted file mode 100644
index 83c14e1183..0000000000
--- a/numba_dpex/dpnp_iface/dpnp_manipulation.py
+++ /dev/null
@@ -1,82 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import numpy as np
-from numba import types
-from numba.core.extending import overload, register_jitable
-from numba.core.typing import signature
-
-import numba_dpex.dpctl_iface as dpctl_functions
-import numba_dpex.dpnp_iface as dpnp_lowering
-import numba_dpex.dpnp_iface.dpnp_stubs_impl as dpnp_ext
-
-from . import stubs
-
-
-@overload(stubs.dpnp.repeat)
-def dpnp_repeat_impl(a, repeats):
-    name = "repeat"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.6.2/dpnp/backend/kernels/dpnp_krnl_manipulation.cpp#L46
-    Function declaration:
-    void dpnp_repeat_c(const void* array1_in, void* result1, const size_t repeats, const size_t size)
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,
-        types.voidptr,
-        types.intp,
-        types.intp,
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a, repeats):
-        if a.size == 0:
-            raise ValueError("Passed Empty array")
-
-        if a.ndim >= 2:
-            raise ValueError("Not supported in dpnp")
-
-        new_size = a.size * repeats
-
-        out = np.zeros(new_size, dtype=a.dtype)
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out_usm = dpctl_functions.malloc_shared(
-            out.size * out.itemsize, sycl_queue
-        )
-
-        dpnp_func(a_usm, out_usm, repeats, a.size)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a.size, out.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-
-        return out
-
-    return dpnp_impl
diff --git a/numba_dpex/dpnp_iface/dpnp_randomimpl.py b/numba_dpex/dpnp_iface/dpnp_randomimpl.py
deleted file mode 100644
index 4274370b48..0000000000
--- a/numba_dpex/dpnp_iface/dpnp_randomimpl.py
+++ /dev/null
@@ -1,1374 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import numpy as np
-from numba import types
-from numba.core.extending import overload, register_jitable
-from numba.core.typing import signature
-
-import numba_dpex.dpctl_iface as dpctl_functions
-import numba_dpex.dpnp_iface as dpnp_lowering
-import numba_dpex.dpnp_iface.dpnp_stubs_impl as dpnp_ext
-
-from . import stubs
-
-
-@register_jitable
-def check_range(low, high):
-    if low >= high:
-        raise ValueError("Low cannot be >= High")
-
-
-@register_jitable
-def common_impl(low, high, res, dpnp_func, print_debug):
-    check_range(low, high)
-
-    sycl_queue = dpctl_functions.get_current_queue()
-    res_usm = dpctl_functions.malloc_shared(res.size * res.itemsize, sycl_queue)
-
-    dpnp_func(res_usm, low, high, res.size)
-
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, res.ctypes, res_usm, res.size * res.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-    dpctl_functions.free_with_queue(res_usm, sycl_queue)
-
-    dpnp_ext._dummy_liveness_func([res.size])
-
-    if print_debug:
-        print("dpnp implementation")
-
-
-@register_jitable
-def common_impl_0_arg(res, dpnp_func, print_debug):
-    sycl_queue = dpctl_functions.get_current_queue()
-    res_usm = dpctl_functions.malloc_shared(res.size * res.itemsize, sycl_queue)
-
-    dpnp_func(res_usm, res.size)
-
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, res.ctypes, res_usm, res.size * res.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-    dpctl_functions.free_with_queue(res_usm, sycl_queue)
-
-    dpnp_ext._dummy_liveness_func([res.size])
-
-    if print_debug:
-        print("dpnp implementation")
-
-
-@register_jitable
-def common_impl_1_arg(arg1, res, dpnp_func, print_debug):
-    sycl_queue = dpctl_functions.get_current_queue()
-    res_usm = dpctl_functions.malloc_shared(res.size * res.itemsize, sycl_queue)
-
-    try:
-        dpnp_func(res_usm, arg1, res.size)
-    except Exception:
-        raise ValueError("Device not supported")
-
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, res.ctypes, res_usm, res.size * res.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-    dpctl_functions.free_with_queue(res_usm, sycl_queue)
-
-    dpnp_ext._dummy_liveness_func([res.size])
-
-    if print_debug:
-        print("dpnp implementation")
-
-
-@register_jitable
-def common_impl_2_arg(arg1, arg2, res, dpnp_func, print_debug):
-    sycl_queue = dpctl_functions.get_current_queue()
-    res_usm = dpctl_functions.malloc_shared(res.size * res.itemsize, sycl_queue)
-
-    dpnp_func(res_usm, arg1, arg2, res.size)
-
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, res.ctypes, res_usm, res.size * res.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-    dpctl_functions.free_with_queue(res_usm, sycl_queue)
-
-    dpnp_ext._dummy_liveness_func([res.size])
-
-    if print_debug:
-        print("dpnp implementation")
-
-
-@register_jitable
-def common_impl_hypergeometric(
-    ngood, nbad, nsample, res, dpnp_func, print_debug
-):
-    sycl_queue = dpctl_functions.get_current_queue()
-    res_usm = dpctl_functions.malloc_shared(res.size * res.itemsize, sycl_queue)
-
-    dpnp_func(res_usm, ngood, nbad, nsample, res.size)
-
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, res.ctypes, res_usm, res.size * res.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    dpctl_functions.free_with_queue(res_usm, sycl_queue)
-
-    dpnp_ext._dummy_liveness_func([res.size])
-
-    if print_debug:
-        print("dpnp implementation")
-
-
-@register_jitable
-def common_impl_multinomial(n, pvals, res, dpnp_func, print_debug):
-    sycl_queue = dpctl_functions.get_current_queue()
-    res_usm = dpctl_functions.malloc_shared(res.size * res.itemsize, sycl_queue)
-
-    pvals_usm = dpctl_functions.malloc_shared(
-        pvals.size * pvals.itemsize, sycl_queue
-    )
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, pvals_usm, pvals.ctypes, pvals.size * pvals.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    dpnp_func(res_usm, n, pvals_usm, pvals.size, res.size)
-
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, res.ctypes, res_usm, res.size * res.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    dpctl_functions.free_with_queue(res_usm, sycl_queue)
-    dpctl_functions.free_with_queue(pvals_usm, sycl_queue)
-
-    dpnp_ext._dummy_liveness_func([res.size])
-
-    if print_debug:
-        print("dpnp implementation")
-
-
-@register_jitable
-def common_impl_multivariate_normal(
-    mean, cov, size, check_valid, tol, res, dpnp_func, print_debug
-):
-    sycl_queue = dpctl_functions.get_current_queue()
-    res_usm = dpctl_functions.malloc_shared(res.size * res.itemsize, sycl_queue)
-
-    mean_usm = dpctl_functions.malloc_shared(
-        mean.size * mean.itemsize, sycl_queue
-    )
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, mean_usm, mean.ctypes, mean.size * mean.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    cov_usm = dpctl_functions.malloc_shared(cov.size * cov.itemsize, sycl_queue)
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, cov_usm, cov.ctypes, cov.size * cov.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    dpnp_func(
-        res_usm, mean.size, mean_usm, mean.size, cov_usm, cov.size, res.size
-    )
-
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, res.ctypes, res_usm, res.size * res.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    dpctl_functions.free_with_queue(res_usm, sycl_queue)
-    dpctl_functions.free_with_queue(mean_usm, sycl_queue)
-    dpctl_functions.free_with_queue(cov_usm, sycl_queue)
-
-    dpnp_ext._dummy_liveness_func([res.size])
-
-    if print_debug:
-        print("dpnp implementation")
-
-
-@overload(stubs.dpnp.random)
-@overload(stubs.dpnp.sample)
-@overload(stubs.dpnp.ranf)
-@overload(stubs.dpnp.random_sample)
-def dpnp_random_impl(size):
-    name = "random_sample"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L391
-
-    Function declaration:
-    void custom_rng_uniform_c(void* result, long low, long high, size_t size)
-
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.int64, types.int64, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-
-    res_dtype = np.float64
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(size):
-        res = np.empty(size, dtype=res_dtype)
-        if res.size != 0:
-            common_impl(0, 1, res, dpnp_func, PRINT_DEBUG)
-        return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.rand)
-def dpnp_random_impl(*size):
-    name = "random_sample"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L391
-
-    Function declaration:
-    void custom_rng_uniform_c(void* result, long low, long high, size_t size)
-
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.int64, types.int64, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-
-    res_dtype = np.float64
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(*size):
-        res = np.empty(size, dtype=res_dtype)
-        if res.size != 0:
-            common_impl(0, 1, res, dpnp_func, PRINT_DEBUG)
-        return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.randint)
-def dpnp_random_impl(low, high=None, size=None):
-    name = "random_sample"
-    dpnp_lowering.ensure_dpnp("randint")
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L391
-
-    Function declaration:
-    void custom_rng_uniform_c(void* result, long low, long high, size_t size)
-
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.int64, types.int64, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["int32", "NONE"], sig)
-
-    res_dtype = np.int32
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if size in (None, types.none):
-        if high not in (None, types.none):
-
-            def dpnp_impl(low, high=None, size=None):
-                res = np.empty(1, dtype=res_dtype)
-                common_impl(low, high, res, dpnp_func, PRINT_DEBUG)
-                return res
-
-        else:
-
-            def dpnp_impl(low, high=None, size=None):
-                res = np.empty(1, dtype=res_dtype)
-                common_impl(0, low, res, dpnp_func, PRINT_DEBUG)
-                return res
-
-    else:
-        if high not in (None, types.none):
-
-            def dpnp_impl(low, high=None, size=None):
-                res = np.empty(size, dtype=res_dtype)
-                if res.size != 0:
-                    common_impl(low, high, res, dpnp_func, PRINT_DEBUG)
-                return res
-
-        else:
-
-            def dpnp_impl(low, high=None, size=None):
-                res = np.empty(size, dtype=res_dtype)
-                if res.size != 0:
-                    common_impl(0, low, res, dpnp_func, PRINT_DEBUG)
-                return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.random_integers)
-def dpnp_random_impl(low, high=None, size=None):
-    name = "random_sample"
-    dpnp_lowering.ensure_dpnp("random_integers")
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L391
-
-    Function declaration:
-    void custom_rng_uniform_c(void* result, long low, long high, size_t size)
-
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.int64, types.int64, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["int32", "NONE"], sig)
-
-    res_dtype = np.int32
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if size in (None, types.none):
-        if high not in (None, types.none):
-
-            def dpnp_impl(low, high=None, size=None):
-                res = np.empty(1, dtype=res_dtype)
-                common_impl(low, high + 1, res, dpnp_func, PRINT_DEBUG)
-                return res
-
-        else:
-
-            def dpnp_impl(low, high=None, size=None):
-                res = np.empty(1, dtype=res_dtype)
-                common_impl(1, low + 1, res, dpnp_func, PRINT_DEBUG)
-                return res
-
-    else:
-        if high not in (None, types.none):
-
-            def dpnp_impl(low, high=None, size=None):
-                res = np.empty(size, dtype=res_dtype)
-                if res.size != 0:
-                    common_impl(low, high + 1, res, dpnp_func, PRINT_DEBUG)
-                return res
-
-        else:
-
-            def dpnp_impl(low, high=None, size=None):
-                res = np.empty(size, dtype=res_dtype)
-                if res.size != 0:
-                    common_impl(1, low + 1, res, dpnp_func, PRINT_DEBUG)
-                return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.beta)
-def dpnp_random_impl(a, b, size=None):
-    name = "beta"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L36
-
-    Function declaration:
-    void custom_rng_beta_c(void* result, _DataType a, _DataType b, size_t size)
-
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.float64, types.float64, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-    res_dtype = np.float64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not (isinstance(a, types.Float)):
-        raise ValueError("We only support float scalar for input: a")
-
-    if not (isinstance(b, types.Float)):
-        raise ValueError("We only support float scalar for input: b")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(a, b, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_2_arg(a, b, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(a, b, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_2_arg(a, b, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.binomial)
-def dpnp_random_impl(n, p, size=None):
-    name = "binomial"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L56
-
-    Function declaration:
-    void custom_rng_binomial_c(void* result, int ntrial, double p, size_t size)
-
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.int32, types.float64, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["int32", "NONE"], sig)
-    res_dtype = np.int32
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not (isinstance(n, types.Integer)):
-        raise ValueError("We only support scalar for input: n")
-
-    if not (isinstance(p, types.Float)):
-        raise ValueError("We only support scalar for input: p")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(n, p, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_2_arg(n, p, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(n, p, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_2_arg(n, p, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.chisquare)
-def dpnp_random_impl(df, size=None):
-    name = "chisquare"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L71
-
-    Function declaration:
-    void custom_rng_chi_square_c(void* result, int df, size_t size)
-    """
-    sig = signature(ret_type, types.voidptr, types.int32, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-    res_dtype = np.float64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not (isinstance(df, types.Integer)):
-        raise ValueError("We only support scalar for input: df")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(df, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_1_arg(df, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(df, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_1_arg(df, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.exponential)
-def dpnp_random_impl(scale=1.0, size=None):
-    name = "exponential"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L86
-
-    Function declaration:
-    void custom_rng_exponential_c(void* result, _DataType beta, size_t size)
-    """
-    sig = signature(ret_type, types.voidptr, types.float64, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-    res_dtype = np.float64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not isinstance(scale, float):
-        if not (isinstance(scale, types.Float)):
-            raise ValueError("We only support scalar for input: scale")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(scale=1.0, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_1_arg(scale, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(scale=1.0, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_1_arg(scale, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.gamma)
-def dpnp_random_impl(shape, scale=1.0, size=None):
-    name = "gamma"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L105
-
-    Function declaration:
-    void custom_rng_gamma_c(void* result, _DataType shape, _DataType scale, size_t size)
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.float64, types.float64, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-    res_dtype = np.float64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not (isinstance(shape, types.Float)):
-        raise ValueError("We only support scalar for input: shape")
-
-    if not isinstance(scale, float):
-        if not (isinstance(scale, types.Float)):
-            raise ValueError("We only support scalar for input: scale")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(shape, scale=1.0, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_2_arg(shape, scale, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(shape, scale=1.0, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_2_arg(shape, scale, res, dpnp_func, PRINT_DEBUG)
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.geometric)
-def dpnp_random_impl(p, size=None):
-    name = "geometric"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L139
-
-    Function declaration:
-    void custom_rng_geometric_c(void* result, float p, size_t size)
-    """
-    sig = signature(ret_type, types.voidptr, types.float32, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["int32", "NONE"], sig)
-    res_dtype = np.int32
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not (isinstance(p, types.Float)):
-        raise ValueError("We only support scalar for input: p")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(p, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_1_arg(p, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(p, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_1_arg(p, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.gumbel)
-def dpnp_random_impl(loc=0.0, scale=1.0, size=None):
-    name = "gumbel"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L154
-
-    Function declaration:
-    void custom_rng_gumbel_c(void* result, double loc, double scale, size_t size)
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.float64, types.float64, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-    res_dtype = np.float64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not isinstance(loc, float):
-        if not (isinstance(loc, types.Float)):
-            raise ValueError("We only support scalar for input: loc")
-
-    if not isinstance(scale, float):
-        if not (isinstance(scale, types.Float)):
-            raise ValueError("We only support scalar for input: scale")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(loc=0.0, scale=1.0, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_2_arg(loc, scale, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(loc=0.0, scale=1.0, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_2_arg(loc, scale, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.hypergeometric)
-def dpnp_random_impl(ngood, nbad, nsample, size=None):
-    name = "hypergeometric"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L169
-
-    Function declaration:
-    void custom_rng_hypergeometric_c(void* result, int l, int s, int m, size_t size)
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,
-        types.int32,
-        types.int32,
-        types.int32,
-        types.intp,
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["int32", "NONE"], sig)
-
-    res_dtype = np.int32
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not (isinstance(ngood, types.Integer)):
-        raise ValueError("We only support scalar for input: ngood")
-
-    if not (isinstance(nbad, types.Integer)):
-        raise ValueError("We only support scalar for input: nbad")
-
-    if not (isinstance(nsample, types.Integer)):
-        raise ValueError("We only support scalar for input: nsample")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(ngood, nbad, nsample, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_hypergeometric(
-                ngood, nbad, nsample, res, dpnp_func, PRINT_DEBUG
-            )
-            return res[0]
-
-    else:
-
-        def dpnp_impl(ngood, nbad, nsample, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_hypergeometric(
-                    ngood, nbad, nsample, res, dpnp_func, PRINT_DEBUG
-                )
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.laplace)
-def dpnp_random_impl(loc=0.0, scale=1.0, size=None):
-    name = "laplace"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L184
-
-    Function declaration:
-    void custom_rng_laplace_c(void* result, double loc, double scale, size_t size)
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.float64, types.float64, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-    res_dtype = np.float64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not isinstance(loc, float):
-        if not (isinstance(loc, types.Float)):
-            raise ValueError("We only support scalar for input: loc")
-
-    if not isinstance(scale, float):
-        if not (isinstance(scale, types.Float)):
-            raise ValueError("We only support scalar for input: scale")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(loc=0.0, scale=1.0, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_2_arg(loc, scale, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(loc=0.0, scale=1.0, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_2_arg(loc, scale, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.lognormal)
-def dpnp_random_impl(mean=0.0, sigma=1.0, size=None):
-    name = "lognormal"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L199
-
-    Function declaration:
-    void custom_rng_lognormal_c(void* result, _DataType mean, _DataType stddev, size_t size)
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.float64, types.float64, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-    res_dtype = np.float64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not isinstance(mean, float):
-        if not (isinstance(mean, types.Float)):
-            raise ValueError("We only support scalar for input: loc")
-
-    if not isinstance(sigma, float):
-        if not (isinstance(sigma, types.Float)):
-            raise ValueError("We only support scalar for input: scale")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(mean=0.0, sigma=1.0, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_2_arg(mean, sigma, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(mean=0.0, sigma=1.0, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_2_arg(mean, sigma, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.multinomial)
-def dpnp_random_impl(n, pvals, size=None):
-    name = "multinomial"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L218
-
-    Function declaration:
-    void custom_rng_multinomial_c(void* result, int ntrial, const double* p_vector,
-                                  const size_t p_vector_size, size_t size)
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,
-        types.int32,
-        types.voidptr,
-        types.intp,
-        types.intp,
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["int32", "NONE"], sig)
-
-    res_dtype = np.int32
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not isinstance(n, types.Integer):
-        raise TypeError(
-            "np.random.multinomial(): n should be an " "integer, got %s" % (n,)
-        )
-
-    if not isinstance(pvals, (types.Sequence, types.Array)):
-        raise TypeError(
-            "np.random.multinomial(): pvals should be an "
-            "array or sequence, got %s" % (pvals,)
-        )
-
-    if size in (None, types.none):
-
-        def dpnp_impl(n, pvals, size=None):
-            out = np.zeros(len(pvals), res_dtype)
-            common_impl_multinomial(n, pvals, out, dpnp_func, PRINT_DEBUG)
-            return out
-
-    elif isinstance(size, types.Integer):
-
-        def dpnp_impl(n, pvals, size=None):
-            out = np.zeros((size, len(pvals)), res_dtype)
-            common_impl_multinomial(n, pvals, out, dpnp_func, PRINT_DEBUG)
-            return out
-
-    elif isinstance(size, types.BaseTuple):
-
-        def dpnp_impl(n, pvals, size=None):
-            out = np.zeros(size + (len(pvals),), res_dtype)
-            common_impl_multinomial(n, pvals, out, dpnp_func, PRINT_DEBUG)
-            return out
-
-    else:
-        raise TypeError(
-            "np.random.multinomial(): size should be int or "
-            "tuple or None, got %s" % (size,)
-        )
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.multivariate_normal)
-def dpnp_random_impl(mean, cov, size=None, check_valid="warn", tol=1e-8):
-    name = "multivariate_normal"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L239
-
-    Function declaration:
-    void custom_rng_multivariate_normal_c(void* result,
-                                      const int dimen,
-                                      const double* mean_vector,
-                                      const size_t mean_vector_size,
-                                      const double* cov_vector,
-                                      const size_t cov_vector_size,
-                                      size_t size)
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,
-        types.int32,
-        types.voidptr,
-        types.intp,
-        types.voidptr,
-        types.intp,
-        types.intp,
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-
-    res_dtype = np.float64
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if size in (None, types.none):
-
-        def dpnp_impl(mean, cov, size=None, check_valid="warn", tol=1e-8):
-            out = np.empty(mean.shape, dtype=res_dtype)
-            common_impl_multivariate_normal(
-                mean, cov, size, check_valid, tol, out, dpnp_func, PRINT_DEBUG
-            )
-            return out
-
-    elif isinstance(size, types.Integer):
-
-        def dpnp_impl(mean, cov, size=None, check_valid="warn", tol=1e-8):
-            new_size = (size,)
-            new_size = new_size + (mean.size,)
-            out = np.empty(new_size, dtype=res_dtype)
-            common_impl_multivariate_normal(
-                mean, cov, size, check_valid, tol, out, dpnp_func, PRINT_DEBUG
-            )
-            return out
-
-    elif isinstance(size, types.BaseTuple):
-
-        def dpnp_impl(mean, cov, size=None, check_valid="warn", tol=1e-8):
-            new_size = size + (mean.size,)
-            out = np.empty(new_size, dtype=res_dtype)
-            common_impl_multivariate_normal(
-                mean, cov, size, check_valid, tol, out, dpnp_func, PRINT_DEBUG
-            )
-            return out
-
-    else:
-        raise TypeError(
-            "np.random.multivariate_normal(): size should be int or "
-            "tuple or None, got %s" % (size,)
-        )
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.negative_binomial)
-def dpnp_random_impl(n, p, size=None):
-    name = "negative_binomial"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L267
-
-    Function declaration:
-    void custom_rng_negative_binomial_c(void* result, double a, double p, size_t size)
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.int32, types.float64, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["int32", "NONE"], sig)
-    res_dtype = np.int32
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not (isinstance(n, types.Integer)):
-        raise ValueError("We only support scalar for input: n")
-
-    if not (isinstance(p, types.Float)):
-        raise ValueError("We only support scalar for input: p")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(n, p, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_2_arg(n, p, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(n, p, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_2_arg(n, p, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.normal)
-def dpnp_random_impl(loc=0.0, scale=1.0, size=None):
-    name = "normal"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L282
-
-    Function declaration:
-    void custom_rng_normal_c(void* result, _DataType mean, _DataType stddev, size_t size)
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.float64, types.float64, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-    res_dtype = np.float64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not isinstance(loc, float):
-        if not (isinstance(loc, types.Float)):
-            raise ValueError("We only support scalar for input: loc")
-
-    if not isinstance(scale, float):
-        if not (isinstance(scale, types.Float)):
-            raise ValueError("We only support scalar for input: scale")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(loc=0.0, scale=1.0, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_2_arg(loc, scale, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(loc=0.0, scale=1.0, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_2_arg(loc, scale, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.poisson)
-def dpnp_random_impl(lam=1.0, size=None):
-    name = "poisson"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L297
-
-    Function declaration:
-    void custom_rng_poisson_c(void* result, double lambda, size_t size)
-    """
-    sig = signature(ret_type, types.voidptr, types.float64, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["int32", "NONE"], sig)
-    res_dtype = np.int32
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not isinstance(lam, float):
-        if not (isinstance(lam, types.Float)):
-            raise ValueError("We only support scalar for input: lam")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(lam=1.0, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_1_arg(lam, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(lam=1.0, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_1_arg(lam, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.rayleigh)
-def dpnp_random_impl(scale=1.0, size=None):
-    name = "rayleigh"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L312
-
-    Function declaration:
-    void custom_rng_rayleigh_c(void* result, _DataType scale, size_t size)
-    """
-    sig = signature(ret_type, types.voidptr, types.float64, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-    res_dtype = np.float64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not isinstance(scale, float):
-        if not (isinstance(scale, types.Float)):
-            raise ValueError("We only support scalar for input: scale")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(scale=1.0, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_1_arg(scale, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(scale=1.0, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_1_arg(scale, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.standard_cauchy)
-def dpnp_random_impl(size=None):
-    name = "standard_cauchy"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L331
-
-    Function declaration:
-    void custom_rng_standard_cauchy_c(void* result, size_t size)
-    """
-    sig = signature(ret_type, types.voidptr, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-    res_dtype = np.float64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if size in (None, types.none):
-
-        def dpnp_impl(size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_0_arg(res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_0_arg(res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.standard_exponential)
-def dpnp_random_impl(size=None):
-    name = "standard_exponential"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L350
-
-    Function declaration:
-    void custom_rng_standard_exponential_c(void* result, size_t size)
-    """
-    sig = signature(ret_type, types.voidptr, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-    res_dtype = np.float64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if size in (None, types.none):
-
-        def dpnp_impl(size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_0_arg(res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_0_arg(res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.standard_gamma)
-def dpnp_random_impl(shape, size=None):
-    name = "standard_gamma"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L364
-
-    Function declaration:
-    void custom_rng_standard_gamma_c(void* result, _DataType shape, size_t size)
-    """
-    sig = signature(ret_type, types.voidptr, types.float64, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-    res_dtype = np.float64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not (isinstance(shape, types.Float)):
-        raise ValueError("We only support scalar for input: shape")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(shape, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_1_arg(shape, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(shape, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_1_arg(shape, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.standard_normal)
-def dpnp_random_impl(size=None):
-    name = "normal"
-    dpnp_lowering.ensure_dpnp("standard_normal")
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L282
-
-    Function declaration:
-    void custom_rng_normal_c(void* result, _DataType mean, _DataType stddev, size_t size)
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.float64, types.float64, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-    res_dtype = np.float64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if size in (None, types.none):
-
-        def dpnp_impl(size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl(0.0, 1.0, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl(0.0, 1.0, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.uniform)
-def dpnp_random_impl(low=0.0, high=1.0, size=None):
-    name = "uniform"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L391
-
-    Function declaration:
-    void custom_rng_uniform_c(void* result, long low, long high, size_t size)
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.int64, types.int64, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-
-    res_dtype = np.float64
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if size in (None, types.none):
-
-        def dpnp_impl(low=0.0, high=1.0, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl(low, high, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    else:
-
-        def dpnp_impl(low=0.0, high=1.0, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl(low, high, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.weibull)
-def dpnp_random_impl(a, size=None):
-    name = "weibull"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_random.cpp#L411
-
-    Function declaration:
-    void custom_rng_weibull_c(void* result, double alpha, size_t size)
-    """
-    sig = signature(ret_type, types.voidptr, types.float64, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, ["float64", "NONE"], sig)
-    res_dtype = np.float64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    if not (isinstance(a, types.Float)):
-        raise ValueError("We only support scalar for input: a")
-
-    if size in (None, types.none):
-
-        def dpnp_impl(a, size=None):
-            res = np.empty(1, dtype=res_dtype)
-            common_impl_1_arg(a, res, dpnp_func, PRINT_DEBUG)
-            return res[0]
-
-    else:
-
-        def dpnp_impl(a, size=None):
-            res = np.empty(size, dtype=res_dtype)
-            if res.size != 0:
-                common_impl_1_arg(a, res, dpnp_func, PRINT_DEBUG)
-            return res
-
-    return dpnp_impl
diff --git a/numba_dpex/dpnp_iface/dpnp_sort_search_countimpl.py b/numba_dpex/dpnp_iface/dpnp_sort_search_countimpl.py
deleted file mode 100644
index 86781d9f5c..0000000000
--- a/numba_dpex/dpnp_iface/dpnp_sort_search_countimpl.py
+++ /dev/null
@@ -1,266 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import numpy as np
-from numba.core import cgutils, types
-from numba.core.extending import overload, register_jitable
-from numba.core.typing import signature
-
-import numba_dpex
-import numba_dpex.dpctl_iface as dpctl_functions
-import numba_dpex.dpnp_iface as dpnp_lowering
-import numba_dpex.dpnp_iface.dpnp_stubs_impl as dpnp_ext
-
-from . import stubs
-
-
-@overload(stubs.dpnp.argmax)
-def dpnp_argmax_impl(a):
-    name = "argmax"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_searching.cpp#L36
-
-    Function declaration:
-    void custom_argmax_c(void* array1_in, void* result1, size_t size)
-    """
-    sig = signature(ret_type, types.voidptr, types.voidptr, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func(
-        "dpnp_" + name, [a.dtype.name, np.dtype(np.int64).name], sig
-    )
-
-    res_dtype = np.int64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        if a.size == 0:
-            raise ValueError("Passed Empty array")
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out = np.empty(1, dtype=res_dtype)
-        out_usm = dpctl_functions.malloc_shared(out.itemsize, sycl_queue)
-
-        dpnp_func(a_usm, out_usm, a.size)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a.size, out.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        return out[0]
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.argmin)
-def dpnp_argmin_impl(a):
-    name = "argmin"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_searching.cpp#L56
-
-    Function declaration:
-    void custom_argmin_c(void* array1_in, void* result1, size_t size)
-    """
-    sig = signature(ret_type, types.voidptr, types.voidptr, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func(
-        "dpnp_" + name, [a.dtype.name, np.dtype(np.int64).name], sig
-    )
-
-    res_dtype = np.int64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        if a.size == 0:
-            raise ValueError("Passed Empty array")
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out = np.empty(1, dtype=res_dtype)
-        out_usm = dpctl_functions.malloc_shared(out.itemsize, sycl_queue)
-
-        dpnp_func(a_usm, out_usm, a.size)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a.size, out.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        return out[0]
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.argsort)
-def dpnp_argsort_impl(a):
-    name = "argsort"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_searching.cpp#L56
-
-    Function declaration:
-    void custom_argmin_c(void* array1_in, void* result1, size_t size)
-    """
-    sig = signature(ret_type, types.voidptr, types.voidptr, types.intp)
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-
-    res_dtype = np.int64
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        if a.size == 0:
-            raise ValueError("Passed Empty array")
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out = np.arange(0, a.size, 1, res_dtype)
-        out_usm = dpctl_functions.malloc_shared(
-            out.size * out.itemsize, sycl_queue
-        )
-
-        dpnp_func(a_usm, out_usm, a.size)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a.size, out.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        return out
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.partition)
-def dpnp_partition_impl(a, kth):
-    name = "partition"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.6.2/dpnp/backend/kernels/dpnp_krnl_sorting.cpp#L90
-    Function declaration:
-    void dpnp_partition_c(
-        void* array1_in, void* array2_in, void* result1, const size_t kth, const size_t* shape_, const size_t ndim)
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,
-        types.voidptr,
-        types.voidptr,
-        types.intp,
-        types.voidptr,
-        types.intp,
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a, kth):
-        if a.size == 0:
-            raise ValueError("Passed Empty array")
-
-        kth_ = kth if kth >= 0 else (a.ndim + kth)
-
-        arr2 = numba_dpex.dpnp.copy(a)
-
-        out = np.empty(a.shape, dtype=a.dtype)
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        arr2_usm = dpctl_functions.malloc_shared(
-            arr2.size * arr2.itemsize, sycl_queue
-        )
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, arr2_usm, arr2.ctypes, arr2.size * arr2.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out_usm = dpctl_functions.malloc_shared(
-            out.size * out.itemsize, sycl_queue
-        )
-
-        dpnp_func(a_usm, arr2_usm, out_usm, kth_, a.shapeptr, a.ndim)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(arr2_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a.size, arr2.size, out.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-
-        return out
-
-    return dpnp_impl
diff --git a/numba_dpex/dpnp_iface/dpnp_statisticsimpl.py b/numba_dpex/dpnp_iface/dpnp_statisticsimpl.py
deleted file mode 100644
index 8adb23eb34..0000000000
--- a/numba_dpex/dpnp_iface/dpnp_statisticsimpl.py
+++ /dev/null
@@ -1,388 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import numpy as np
-from numba.core import cgutils, types
-from numba.core.extending import overload, register_jitable
-from numba.core.typing import signature
-
-import numba_dpex.dpctl_iface as dpctl_functions
-import numba_dpex.dpnp_iface as dpnp_lowering
-import numba_dpex.dpnp_iface.dpnp_stubs_impl as dpnp_ext
-
-from . import stubs
-
-
-@overload(stubs.dpnp.max)
-@overload(stubs.dpnp.amax)
-def dpnp_amax_impl(a):
-    name = "max"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/e389248c709531b181be8bf33b1a270fca812a92/dpnp/backend/kernels/dpnp_krnl_statistics.cpp#L149
-
-    Function declaration:
-    void dpnp_max_c(void* array1_in, void* result1, const size_t result_size, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis)
-
-    We are using void * in case of size_t * as Numba currently does not have
-    any type to represent size_t *. Since, both the types are pointers,
-    if the compiler allows there should not be any mismatch in the size of
-    the container to hold different types of pointer.
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,
-        types.voidptr,
-        types.intp,
-        types.voidptr,
-        types.intp,
-        types.voidptr,
-        types.intp,
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        if a.size == 0:
-            raise ValueError("Passed Empty array")
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out_usm = dpctl_functions.malloc_shared(a.itemsize, sycl_queue)
-
-        axis, naxis = 0, 0
-
-        dpnp_func(
-            a_usm, out_usm, a.size * a.itemsize, a.shapeptr, a.ndim, axis, naxis
-        )
-
-        out = np.empty(1, dtype=a.dtype)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a.size, out.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        return out[0]
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.min)
-@overload(stubs.dpnp.amin)
-def dpnp_amin_impl(a):
-    name = "min"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/57caae8beb607992f40cdbe00f2666ee84358a97/dpnp/backend/kernels/dpnp_krnl_statistics.cpp#L412
-
-    Function declaration:
-    void dpnp_min_c(void* array1_in, void* result1, const size_t result_size, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis)
-
-    We are using void * in case of size_t * as Numba currently does not have
-    any type to represent size_t *. Since, both the types are pointers,
-    if the compiler allows there should not be any mismatch in the size of
-    the container to hold different types of pointer.
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,
-        types.voidptr,
-        types.intp,
-        types.voidptr,
-        types.intp,
-        types.voidptr,
-        types.intp,
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        if a.size == 0:
-            raise ValueError("Passed Empty array")
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out_usm = dpctl_functions.malloc_shared(a.itemsize, sycl_queue)
-
-        dpnp_func(
-            a_usm,
-            out_usm,
-            a.size * a.itemsize,
-            a.shapeptr,
-            a.ndim,
-            a.shapeptr,
-            0,
-        )
-
-        out = np.empty(1, dtype=a.dtype)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a.size, out.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        return out[0]
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.mean)
-def dpnp_mean_impl(a):
-    name = "mean"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.6.1dev/dpnp/backend/kernels/dpnp_krnl_statistics.cpp#L185
-
-    Function declaration:
-    void dpnp_mean_c(void* array1_in, void* result1, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis)
-
-    We are using void * in case of size_t * as Numba currently does not have
-    any type to represent size_t *. Since, both the types are pointers,
-    if the compiler allows there should not be any mismatch in the size of
-    the container to hold different types of pointer.
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,
-        types.voidptr,
-        types.voidptr,
-        types.intp,
-        types.voidptr,
-        types.intp,
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    res_dtype = np.float64
-    if a.dtype == types.float32:
-        res_dtype = np.float32
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        if a.size == 0:
-            raise ValueError("Passed Empty array")
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out = np.empty(1, dtype=res_dtype)
-        out_usm = dpctl_functions.malloc_shared(out.itemsize, sycl_queue)
-
-        axis, naxis = 0, 0
-
-        dpnp_func(a_usm, out_usm, a.shapeptr, a.ndim, axis, naxis)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a.size, out.size])
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        return out[0]
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.median)
-def dpnp_median_impl(a):
-    name = "median"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_statistics.cpp#L213
-
-    Function declaration:
-    void custom_median_c(void* array1_in, void* result1, const size_t* shape,
-                         size_t ndim, const size_t* axis, size_t naxis)
-
-    We are using void * in case of size_t * as Numba currently does not have
-    any type to represent size_t *. Since, both the types are pointers,
-    if the compiler allows there should not be any mismatch in the size of
-    the container to hold different types of pointer.
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,
-        types.voidptr,
-        types.voidptr,
-        types.intp,
-        types.voidptr,
-        types.intp,
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-
-    res_dtype = np.float64
-    if a.dtype == types.float32:
-        res_dtype = np.float32
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        if a.size == 0:
-            raise ValueError("Passed Empty array")
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        out = np.empty(1, dtype=res_dtype)
-        out_usm = dpctl_functions.malloc_shared(out.itemsize, sycl_queue)
-
-        dpnp_func(a_usm, out_usm, a.shapeptr, a.ndim, a.shapeptr, a.ndim)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a.size, out.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        return out[0]
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.cov)
-def dpnp_cov_impl(a):
-    name = "cov"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.4.0/dpnp/backend/custom_kernels_statistics.cpp#L51
-
-    Function declaration:
-    void custom_cov_c(void* array1_in, void* result1, size_t nrows, size_t ncols)
-    """
-    sig = signature(
-        ret_type, types.voidptr, types.voidptr, types.intp, types.intp
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-
-    res_dtype = np.float64
-    copy_input_to_double = True
-    if a.dtype == types.float64:
-        copy_input_to_double = False
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        if a.size == 0:
-            raise ValueError("Passed Empty array")
-
-        sycl_queue = dpctl_functions.get_current_queue()
-
-        """ We have to pass a array in double precision to DpNp """
-        if copy_input_to_double:
-            a_copy_in_double = a.astype(np.float64)
-        else:
-            a_copy_in_double = a
-        a_usm = dpctl_functions.malloc_shared(
-            a_copy_in_double.size * a_copy_in_double.itemsize, sycl_queue
-        )
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue,
-            a_usm,
-            a_copy_in_double.ctypes,
-            a_copy_in_double.size * a_copy_in_double.itemsize,
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        if a.ndim == 2:
-            rows = a.shape[0]
-            cols = a.shape[1]
-            out = np.empty((rows, rows), dtype=res_dtype)
-        elif a.ndim == 1:
-            rows = 1
-            cols = a.shape[0]
-            out = np.empty(rows, dtype=res_dtype)
-
-        out_usm = dpctl_functions.malloc_shared(
-            out.size * out.itemsize, sycl_queue
-        )
-
-        dpnp_func(a_usm, out_usm, rows, cols)
-
-        event = dpctl_functions.queue_memcpy(
-            sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-        )
-        dpctl_functions.event_wait(event)
-        dpctl_functions.event_delete(event)
-
-        dpctl_functions.free_with_queue(a_usm, sycl_queue)
-        dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-        dpnp_ext._dummy_liveness_func([a_copy_in_double.size, a.size, out.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-        if a.ndim == 2:
-            return out
-        elif a.ndim == 1:
-            return out[0]
-
-    return dpnp_impl
diff --git a/numba_dpex/dpnp_iface/dpnp_stubs_impl.py b/numba_dpex/dpnp_iface/dpnp_stubs_impl.py
deleted file mode 100644
index e1317dda03..0000000000
--- a/numba_dpex/dpnp_iface/dpnp_stubs_impl.py
+++ /dev/null
@@ -1,79 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import numpy as np
-from llvmlite import ir
-from numba.core import types
-from numba.core.extending import register_jitable
-from numba.core.imputils import lower_getattr
-from numba.cpython import listobj
-
-ll_void_p = ir.IntType(8).as_pointer()
-
-
-def get_dpnp_fptr(fn_name, type_names):
-    from . import dpnp_fptr_interface as dpnp_iface
-
-    f_ptr = dpnp_iface.get_dpnp_fn_ptr(fn_name, type_names)
-    return f_ptr
-
-
-@register_jitable
-def _check_finite_matrix(a):
-    for v in np.nditer(a):
-        if not np.isfinite(v.item()):
-            raise np.linalg.LinAlgError("Array must not contain infs or NaNs.")
-
-
-@register_jitable
-def _dummy_liveness_func(a):
-    """pass a list of variables to be preserved through dead code elimination"""
-    return a[0]
-
-
-def dpnp_func(fn_name, type_names, sig):
-    f_ptr = get_dpnp_fptr(fn_name, type_names)
-
-    def get_pointer(obj):
-        return f_ptr
-
-    return types.ExternalFunctionPointer(sig, get_pointer=get_pointer)
-
-
-"""
-This function retrieves the pointer to the structure where the shape
-of an ndarray is stored. We cast it to void * to make it easier to
-pass around.
-"""
-
-
-@lower_getattr(types.Array, "shapeptr")
-def array_shapeptr(context, builder, typ, value):
-    shape_ptr = builder.gep(
-        value.operands[0],
-        [
-            context.get_constant(types.int32, 0),
-            context.get_constant(types.int32, 5),
-        ],
-    )
-
-    return builder.bitcast(shape_ptr, ll_void_p)
-
-
-@lower_getattr(types.List, "size")
-def list_size(context, builder, typ, value):
-    inst = listobj.ListInstance(context, builder, typ, value)
-    return inst.size
-
-
-@lower_getattr(types.List, "itemsize")
-def list_itemsize(context, builder, typ, value):
-    llty = context.get_data_type(typ.dtype)
-    return context.get_constant(types.uintp, context.get_abi_sizeof(llty))
-
-
-@lower_getattr(types.List, "ctypes")
-def list_ctypes(context, builder, typ, value):
-    inst = listobj.ListInstance(context, builder, typ, value)
-    return builder.bitcast(inst.data, ll_void_p)
diff --git a/numba_dpex/dpnp_iface/dpnp_transcendentalsimpl.py b/numba_dpex/dpnp_iface/dpnp_transcendentalsimpl.py
deleted file mode 100644
index fd76b1b83b..0000000000
--- a/numba_dpex/dpnp_iface/dpnp_transcendentalsimpl.py
+++ /dev/null
@@ -1,194 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import numpy as np
-from numba import types
-from numba.core.extending import overload, register_jitable
-from numba.core.typing import signature
-
-import numba_dpex
-import numba_dpex.dpctl_iface as dpctl_functions
-import numba_dpex.dpnp_iface as dpnp_lowering
-import numba_dpex.dpnp_iface.dpnp_stubs_impl as dpnp_ext
-
-from . import stubs
-
-
-@register_jitable
-def common_impl(a, out, dpnp_func, print_debug):
-    if a.size == 0:
-        raise ValueError("Passed Empty array")
-
-    sycl_queue = dpctl_functions.get_current_queue()
-    a_usm = dpctl_functions.malloc_shared(a.size * a.itemsize, sycl_queue)
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, a_usm, a.ctypes, a.size * a.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    out_usm = dpctl_functions.malloc_shared(a.itemsize, sycl_queue)
-
-    axes, axes_ndim = 0, 0
-    initial = 0
-    where = 0
-
-    dpnp_func(
-        out_usm, a_usm, a.shapeptr, a.ndim, axes, axes_ndim, initial, where
-    )
-
-    event = dpctl_functions.queue_memcpy(
-        sycl_queue, out.ctypes, out_usm, out.size * out.itemsize
-    )
-    dpctl_functions.event_wait(event)
-    dpctl_functions.event_delete(event)
-
-    dpctl_functions.free_with_queue(a_usm, sycl_queue)
-    dpctl_functions.free_with_queue(out_usm, sycl_queue)
-
-    dpnp_ext._dummy_liveness_func([a.size, out.size])
-
-    if print_debug:
-        print("dpnp implementation")
-
-
-@overload(stubs.dpnp.sum)
-def dpnp_sum_impl(a):
-    name = "sum"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.6.1dev/dpnp/backend/kernels/dpnp_krnl_reduction.cpp#L59
-
-    Function declaration:
-    void dpnp_sum_c(void* result_out,
-                    const void* input_in,
-                    const size_t* input_shape,
-                    const size_t input_shape_ndim,
-                    const long* axes,
-                    const size_t axes_ndim,
-                    const void* initial,
-                    const long* where)
-
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,  # void* result_out,
-        types.voidptr,  # const void* input_in,
-        types.voidptr,  # const size_t* input_shape,
-        types.intp,  # const size_t input_shape_ndim,
-        types.voidptr,  # const long* axes,
-        types.intp,  # const size_t axes_ndim,
-        types.voidptr,  # const void* initial,
-        types.voidptr,  # const long* where)
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        out = np.empty(1, dtype=a.dtype)
-        common_impl(a, out, dpnp_func, PRINT_DEBUG)
-
-        return out[0]
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.prod)
-def dpnp_prod_impl(a):
-    name = "prod"
-    dpnp_lowering.ensure_dpnp(name)
-
-    ret_type = types.void
-    """
-    dpnp source:
-    https://github.com/IntelPython/dpnp/blob/0.6.1dev/dpnp/backend/kernels/dpnp_krnl_reduction.cpp#L129
-
-    Function declaration:
-    void dpnp_prod_c(void* result_out,
-                     const void* input_in,
-                     const size_t* input_shape,
-                     const size_t input_shape_ndim,
-                     const long* axes,
-                     const size_t axes_ndim,
-                     const void* initial, // type must be _DataType_output
-                     const long* where)
-    """
-    sig = signature(
-        ret_type,
-        types.voidptr,  # void* result_out,
-        types.voidptr,  # const void* input_in,
-        types.voidptr,  # const size_t* input_shape,
-        types.intp,  # const size_t input_shape_ndim,
-        types.voidptr,  # const long* axes,
-        types.intp,  # const long* axes,
-        types.voidptr,  # const void* initial, // type must be _DataType_output
-        types.voidptr,  # const long* where)
-    )
-    dpnp_func = dpnp_ext.dpnp_func("dpnp_" + name, [a.dtype.name, "NONE"], sig)
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        out = np.empty(1, dtype=a.dtype)
-
-        common_impl(a, out, dpnp_func, PRINT_DEBUG)
-        return out[0]
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.nansum)
-def dpnp_nansum_impl(a):
-    name = "nansum"
-    dpnp_lowering.ensure_dpnp(name)
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        a_ravel = a.ravel()
-        a_ravel_copy = np.copy(a_ravel)
-
-        for i in range(len(a_ravel_copy)):
-            if np.isnan(a_ravel_copy[i]):
-                a_ravel_copy[i] = 0
-
-        result = numba_dpex.dpnp.sum(a_ravel_copy)
-        dpnp_ext._dummy_liveness_func([a.size, a_ravel_copy.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-
-        return result
-
-    return dpnp_impl
-
-
-@overload(stubs.dpnp.nanprod)
-def dpnp_nanprod_impl(a):
-    name = "nanprod"
-    dpnp_lowering.ensure_dpnp(name)
-
-    PRINT_DEBUG = dpnp_lowering.DEBUG
-
-    def dpnp_impl(a):
-        a_ravel = a.ravel()
-        a_ravel_copy = np.copy(a_ravel)
-
-        for i in range(len(a_ravel_copy)):
-            if np.isnan(a_ravel_copy[i]):
-                a_ravel_copy[i] = 1
-
-        result = numba_dpex.dpnp.prod(a_ravel_copy)
-        dpnp_ext._dummy_liveness_func([a.size, a_ravel_copy.size])
-
-        if PRINT_DEBUG:
-            print("dpnp implementation")
-
-        return result
-
-    return dpnp_impl
diff --git a/numba_dpex/dpnp_iface/dpnpdecl.py b/numba_dpex/dpnp_iface/dpnpdecl.py
deleted file mode 100644
index 8ee74a0ca3..0000000000
--- a/numba_dpex/dpnp_iface/dpnpdecl.py
+++ /dev/null
@@ -1,45 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-from numba import types
-from numba.core.typing.templates import AttributeTemplate, infer_getattr
-
-import numba_dpex
-
-
-@infer_getattr
-class DpnpTemplate(AttributeTemplate):
-    key = types.Module(numba_dpex)
-
-    def resolve_dpnp(self, mod):
-        return types.Module(numba_dpex.dpnp)
-
-
-"""
-This adds a shapeptr attribute to Numba type representing np.ndarray.
-This allows us to get the raw pointer to the structure where the shape
-of an ndarray is stored from an overloaded implementation
-"""
-
-
-@infer_getattr
-class ArrayAttribute(AttributeTemplate):
-    key = types.Array
-
-    def resolve_shapeptr(self, ary):
-        return types.voidptr
-
-
-@infer_getattr
-class ListAttribute(AttributeTemplate):
-    key = types.List
-
-    def resolve_size(self, ary):
-        return types.int64
-
-    def resolve_itemsize(self, ary):
-        return types.int64
-
-    def resolve_ctypes(self, ary):
-        return types.voidptr
diff --git a/numba_dpex/dpnp_iface/stubs.py b/numba_dpex/dpnp_iface/stubs.py
deleted file mode 100644
index 0c7708dca6..0000000000
--- a/numba_dpex/dpnp_iface/stubs.py
+++ /dev/null
@@ -1,197 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-from numba_dpex.ocl.stubs import Stub
-
-
-class dpnp(Stub):
-    """dpnp namespace"""
-
-    _description_ = "<dpnp>"
-
-    class all(Stub):
-        pass
-
-    class amax(Stub):
-        pass
-
-    class amin(Stub):
-        pass
-
-    class argmax(Stub):
-        pass
-
-    class argmin(Stub):
-        pass
-
-    class argsort(Stub):
-        pass
-
-    class beta(Stub):
-        pass
-
-    class binomial(Stub):
-        pass
-
-    class chisquare(Stub):
-        pass
-
-    class cholesky(Stub):
-        pass
-
-    class copy(Stub):
-        pass
-
-    class cov(Stub):
-        pass
-
-    class cumprod(Stub):
-        pass
-
-    class cumsum(Stub):
-        pass
-
-    class det(Stub):
-        pass
-
-    class diagonal(Stub):
-        pass
-
-    class dot(Stub):
-        pass
-
-    class eig(Stub):
-        pass
-
-    class eigvals(Stub):
-        pass
-
-    class exponential(Stub):
-        pass
-
-    class gamma(Stub):
-        pass
-
-    class geometric(Stub):
-        pass
-
-    class gumbel(Stub):
-        pass
-
-    class hypergeometric(Stub):
-        pass
-
-    class laplace(Stub):
-        pass
-
-    class lognormal(Stub):
-        pass
-
-    class matmul(Stub):
-        pass
-
-    class matrix_power(Stub):
-        pass
-
-    class matrix_rank(Stub):
-        pass
-
-    class max(Stub):
-        pass
-
-    class mean(Stub):
-        pass
-
-    class median(Stub):
-        pass
-
-    class min(Stub):
-        pass
-
-    class multi_dot(Stub):
-        pass
-
-    class multinomial(Stub):
-        pass
-
-    class multivariate_normal(Stub):
-        pass
-
-    class nanprod(Stub):
-        pass
-
-    class nansum(Stub):
-        pass
-
-    class negative_binomial(Stub):
-        pass
-
-    class normal(Stub):
-        pass
-
-    class partition(Stub):
-        pass
-
-    class poisson(Stub):
-        pass
-
-    class prod(Stub):
-        pass
-
-    class rand(Stub):
-        pass
-
-    class randint(Stub):
-        pass
-
-    class random_integers(Stub):
-        pass
-
-    class random_sample(Stub):
-        pass
-
-    class random(Stub):
-        pass
-
-    class ranf(Stub):
-        pass
-
-    class rayleigh(Stub):
-        pass
-
-    class repeat(Stub):
-        pass
-
-    class sample(Stub):
-        pass
-
-    class sort(Stub):
-        pass
-
-    class standard_cauchy(Stub):
-        pass
-
-    class standard_exponential(Stub):
-        pass
-
-    class standard_gamma(Stub):
-        pass
-
-    class standard_normal(Stub):
-        pass
-
-    class sum(Stub):
-        pass
-
-    class take(Stub):
-        pass
-
-    class uniform(Stub):
-        pass
-
-    class vdot(Stub):
-        pass
-
-    class weibull(Stub):
-        pass
diff --git a/numba_dpex/examples/kernel/select_device_for_kernel.py b/numba_dpex/examples/kernel/select_device_for_kernel.py
deleted file mode 100644
index df19d0f02d..0000000000
--- a/numba_dpex/examples/kernel/select_device_for_kernel.py
+++ /dev/null
@@ -1,127 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import warnings
-
-import dpctl
-import dpctl.tensor as dpt
-import numpy as np
-
-import numba_dpex as ndpx
-
-"""
-We support passing arrays of two types to a @numba_dpex.kernel decorated
-function.
-
-1. numpy.ndarray
-2. Any array with __sycl_usm_array_interface__ (SUAI) attribute.
-
-Users are not allowed to mix the type of arrays passed as arguments. As in, all
-the arguments passed to a @numba_dpex.kernel has to have the same type. For
-example, if the first array argument is of type numpy.ndarray the rest of the
-array arguments will also have to be of type numpy.ndarray.
-
-The following are how users can specify in which device they want to offload
-their computation.
-
-1. numpy.ndarray
-    Using context manager provided by ndpx. Please look at method:
-        select_device_ndarray()
-
-2. Array with __sycl_usm_array_interface__ attribute
-    We follow compute follows data which states that the device where the
-    data resides will be selected for computation. Please look at method:
-         select_device_SUAI()
-
-    Users can mix SUAI arrays created using equivalent SYCL queues.
-    Two SYCL queues are equivalent if they have the same:
-        1. SYCL context
-        2. SYCL device
-        3. Same queue properties
-"""
-
-
-@ndpx.kernel
-def sum_kernel(a, b, c):
-    i = ndpx.get_global_id(0)
-    c[i] = a[i] + b[i]
-
-
-def allocate_SUAI_data(a, b, got, usm_type, queue):
-    da = dpt.usm_ndarray(
-        a.shape,
-        dtype=a.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue},
-    )
-    da.usm_data.copy_from_host(a.reshape((-1)).view("|u1"))
-
-    db = dpt.usm_ndarray(
-        b.shape,
-        dtype=b.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue},
-    )
-    db.usm_data.copy_from_host(b.reshape((-1)).view("|u1"))
-
-    dc = dpt.usm_ndarray(
-        got.shape,
-        dtype=got.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue},
-    )
-
-    return da, db, dc
-
-
-# ==========================================================================
-def select_device_ndarray(N):
-    a = np.array(np.random.random(N), np.float32)
-    b = np.array(np.random.random(N), np.float32)
-
-    got = np.ones_like(a)
-
-    # This context manager is specifying to use the Opencl GPU.
-    default_device = dpctl.select_default_device()
-
-    with ndpx.offload_to_sycl_device(default_device.filter_string):
-        sum_kernel[ndpx.NdRange(ndpx.Range(N), ndpx.Range(1))](a, b, got)
-
-    expected = a + b
-
-    assert np.array_equal(got, expected)
-    print("Correct result when numpy.ndarray is passed!")
-
-
-def select_device_SUAI(N):
-    usm_type = "device"
-
-    a = np.array(np.random.random(N), np.float32)
-    b = np.array(np.random.random(N), np.float32)
-    got = np.ones_like(a)
-
-    device = dpctl.select_default_device()
-    queue = dpctl.SyclQueue(device)
-
-    # We are allocating the data in Opencl GPU and this device
-    # will be selected for compute.
-    da, db, dc = allocate_SUAI_data(a, b, got, usm_type, queue)
-
-    # Users don't need to specify where the computation will
-    # take place. It will be inferred from data.
-    sum_kernel[ndpx.NdRange(ndpx.Range(N), ndpx.Range(1))](da, db, dc)
-
-    dc.usm_data.copy_to_host(got.reshape((-1)).view("|u1"))
-
-    expected = a + b
-
-    assert np.array_equal(got, expected)
-    print(
-        "Correct result when array with __sycl_usm_array_interface__ is passed!"
-    )
-
-
-if __name__ == "__main__":
-    select_device_ndarray(10)
-    select_device_SUAI(10)
diff --git a/numba_dpex/retarget.py b/numba_dpex/retarget.py
deleted file mode 100644
index b327dd6cb3..0000000000
--- a/numba_dpex/retarget.py
+++ /dev/null
@@ -1,61 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-from contextlib import contextmanager
-
-import dpctl
-from numba import njit
-
-try:
-    from numba.core.dispatcher import TargetConfigurationStack
-except ImportError:
-    # for support numba 0.54 and <=0.55.0dev0=*_469
-    from numba.core.dispatcher import TargetConfig as TargetConfigurationStack
-
-from numba.core.retarget import BasicRetarget
-
-from numba_dpex.core.targets.kernel_target import DPEX_KERNEL_TARGET_NAME
-
-
-class DpexRetarget(BasicRetarget):
-    def __init__(self, filter_str):
-        self.filter_str = filter_str
-        super(DpexRetarget, self).__init__()
-
-    @property
-    def output_target(self):
-        return DPEX_KERNEL_TARGET_NAME
-
-    def compile_retarget(self, cpu_disp):
-        kernel = njit(_target=DPEX_KERNEL_TARGET_NAME)(cpu_disp.py_func)
-        return kernel
-
-
-_first_level_cache = dict()
-
-
-def _retarget(sycl_queue):
-    filter_string = sycl_queue.sycl_device.filter_string
-
-    result = _first_level_cache.get(filter_string)
-
-    if not result:
-        result = DpexRetarget(filter_string)
-        _first_level_cache[filter_string] = result
-
-    return result
-
-
-def _retarget_context_manager(sycl_queue):
-    """Return context manager for retargeting njit offloading."""
-    retarget = _retarget(sycl_queue)
-    return TargetConfigurationStack.switch_target(retarget)
-
-
-def _register_context_factory():
-    dpctl.nested_context_factories.append(_retarget_context_manager)
-
-
-_register_context_factory()
-offload_to_sycl_device = dpctl.device_context
diff --git a/setup.py b/setup.py
index e571c4331b..f1eebd7309 100644
--- a/setup.py
+++ b/setup.py
@@ -45,20 +45,6 @@ def get_ext_modules():
     if IS_LIN:
         dpctl_runtime_library_dirs.append(os.path.dirname(dpctl.__file__))
 
-    if dpnp_present:
-        dpnp_lib_path = []
-        dpnp_lib_path += [os.path.dirname(dpnp.__file__)]
-        ext_dpnp_iface = Extension(
-            name="numba_dpex.dpnp_iface.dpnp_fptr_interface",
-            sources=["numba_dpex/dpnp_iface/dpnp_fptr_interface.pyx"],
-            include_dirs=[dpnp.get_include(), dpctl.get_include()],
-            libraries=["dpnp_backend_c"],
-            library_dirs=dpnp_lib_path,
-            runtime_library_dirs=(dpnp_lib_path if IS_LIN else []),
-            language="c++",
-        )
-        ext_modules += [ext_dpnp_iface]
-
     ext_dpexrt_python = Extension(
         name="numba_dpex.core.runtime._dpexrt_python",
         sources=[


Test patch:
diff --git a/numba_dpex/tests/_helper.py b/numba_dpex/tests/_helper.py
index 515a4a31d8..c3c6b32ee7 100644
--- a/numba_dpex/tests/_helper.py
+++ b/numba_dpex/tests/_helper.py
@@ -9,7 +9,6 @@
 
 import dpctl
 import pytest
-from numba.tests.support import captured_stdout
 
 from numba_dpex import config, numba_version
 
@@ -139,76 +138,3 @@ def override_config(name, value, config=config):
 
 def _id(obj):
     return obj
-
-
-def _ensure_dpnp():
-    try:
-        from numba_dpex.dpnp_iface import dpnp_fptr_interface as dpnp_iface
-
-        return True
-    except ImportError:
-        if config.TESTING_SKIP_NO_DPNP:
-            return False
-        else:
-            pytest.fail("DPNP is not available")
-
-
-skip_no_dpnp = pytest.mark.skipif(
-    not _ensure_dpnp(), reason="DPNP is not available"
-)
-
-
-@contextlib.contextmanager
-def dpnp_debug():
-    import numba_dpex.dpnp_iface as dpnp_lowering
-
-    old, dpnp_lowering.DEBUG = dpnp_lowering.DEBUG, 1
-    yield
-    dpnp_lowering.DEBUG = old
-
-
-@contextlib.contextmanager
-def assert_dpnp_implementaion():
-    from numba.tests.support import captured_stdout
-
-    with captured_stdout() as stdout, dpnp_debug():
-        yield
-
-    assert (
-        "dpnp implementation" in stdout.getvalue()
-    ), "dpnp implementation is not used"
-
-
-@contextlib.contextmanager
-def assert_auto_offloading(parfor_offloaded=1, parfor_offloaded_failure=0):
-    """
-    If ``parfor_offloaded`` is not provided this context_manager
-    will check for 1 occurrance of success message. Developers
-    can always specify how many parfor offload success message
-    is expected.
-    If ``parfor_offloaded_failure`` is not provided the default
-    behavior is to expect 0 failure message, in other words, we
-    expect all parfors present in the code to be successfully
-    offloaded to GPU.
-    """
-    old_debug = config.DEBUG
-    config.DEBUG = 1
-
-    with captured_stdout() as stdout:
-        yield
-
-    config.DEBUG = old_debug
-
-    got_parfor_offloaded = stdout.getvalue().count("Parfor offloaded to")
-    assert parfor_offloaded == got_parfor_offloaded, (
-        "Expected %d parfor(s) to be auto offloaded, instead got %d parfor(s) auto offloaded"
-        % (parfor_offloaded, got_parfor_offloaded)
-    )
-
-    got_parfor_offloaded_failure = stdout.getvalue().count(
-        "Failed to offload parfor to"
-    )
-    assert parfor_offloaded_failure == got_parfor_offloaded_failure, (
-        "Expected %d parfor(s) to be not auto offloaded, instead got %d parfor(s) not auto offloaded"
-        % (parfor_offloaded_failure, got_parfor_offloaded_failure)
-    )
diff --git a/numba_dpex/tests/core/passes/test_rename_numpy_function_pass.py b/numba_dpex/tests/core/passes/test_rename_numpy_function_pass.py
deleted file mode 100644
index 7cb539e39a..0000000000
--- a/numba_dpex/tests/core/passes/test_rename_numpy_function_pass.py
+++ /dev/null
@@ -1,113 +0,0 @@
-#! /usr/bin/env python
-
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import numba
-import numpy as np
-import pytest
-from numba import typeof
-from numba.core import compiler, cpu, typing
-from numba.core.typed_passes import AnnotateTypes, NopythonTypeInference
-
-import numba_dpex
-from numba_dpex.core.passes.rename_numpy_functions_pass import (
-    RewriteNdarrayFunctionsPass,
-    RewriteOverloadedNumPyFunctionsPass,
-)
-from numba_dpex.tests._helper import skip_no_dpnp
-
-
-class MyPipeline(object):
-    def __init__(self, test_ir, args):
-        self.state = compiler.StateDict()
-        self.state.typingctx = typing.Context()
-        self.state.targetctx = cpu.CPUContext(self.state.typingctx)
-        self.state.func_ir = test_ir
-        self.state.func_id = test_ir.func_id
-        self.state.args = args
-        self.state.return_type = None
-        self.state.locals = dict()
-        self.state.status = None
-        self.state.lifted = dict()
-        self.state.lifted_from = None
-        self.state.typingctx.refresh()
-        self.state.targetctx.refresh()
-
-
-def check_equivalent(expected_ir, got_ir):
-    expected_block_body = expected_ir.blocks[0].body
-    got_block_body = got_ir.blocks[0].body
-
-    if len(expected_block_body) != len(got_block_body):
-        return False
-
-    for i in range(len(expected_block_body)):
-        expected_stmt = expected_block_body[i]
-        got_stmt = got_block_body[i]
-        if type(expected_stmt) != type(got_stmt):
-            return False
-        else:
-            if isinstance(expected_stmt, numba.core.ir.Assign):
-                if isinstance(expected_stmt.value, numba.core.ir.Global):
-                    if (
-                        expected_stmt.value.name != got_stmt.value.name
-                        and expected_stmt.value.name != "numba_dpex"
-                    ):
-                        return False
-                elif isinstance(expected_stmt.value, numba.core.ir.Expr):
-                    # should get "dpnp" and "sum" as attr
-                    if expected_stmt.value.op == "getattr":
-                        if expected_stmt.value.attr != got_stmt.value.attr:
-                            return False
-    return True
-
-
-class TestRenameNumpyFunctionsPass:
-    def test_rename_numpy(self):
-        def expected(a):
-            return numba_dpex.dpnp.sum(a)
-
-        def got(a):
-            return np.sum(a)
-
-        expected_ir = compiler.run_frontend(expected)
-        got_ir = compiler.run_frontend(got)
-
-        pipeline = MyPipeline(got_ir, None)
-
-        rewrite_numpy_functions_pass = RewriteOverloadedNumPyFunctionsPass()
-        rewrite_numpy_functions_pass.run_pass(pipeline.state)
-
-        assert check_equivalent(expected_ir, pipeline.state.func_ir)
-
-
-@skip_no_dpnp
-class TestRenameNdarrayFunctionsPass:
-    def test_rename_ndarray(self):
-        def expected(a):
-            return numba_dpex.dpnp.sum(a)
-
-        def got(a):
-            return a.sum()
-
-        expected_ir = compiler.run_frontend(expected)
-        got_ir = compiler.run_frontend(got)
-
-        a = np.arange(10)
-        args = [a]
-        argtypes = [typeof(x) for x in args]
-
-        pipeline = MyPipeline(got_ir, argtypes)
-
-        tyinfer_pass = NopythonTypeInference()
-        tyinfer_pass.run_pass(pipeline.state)
-
-        annotate_ty_pass = AnnotateTypes()
-        annotate_ty_pass.run_pass(pipeline.state)
-
-        rewrite_ndarray_functions_pass = RewriteNdarrayFunctionsPass()
-        rewrite_ndarray_functions_pass.run_pass(pipeline.state)
-
-        assert check_equivalent(expected_ir, pipeline.state.func_ir)
diff --git a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_bitwise_ops.py b/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_bitwise_ops.py
index 7662df8afc..e5ecf7cbbf 100644
--- a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_bitwise_ops.py
+++ b/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_bitwise_ops.py
@@ -9,7 +9,7 @@
 import pytest
 
 from numba_dpex import dpjit
-from numba_dpex.tests._helper import assert_auto_offloading, filter_strings
+from numba_dpex.tests._helper import filter_strings
 
 list_of_binary_ops = [
     "bitwise_and",
diff --git a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_logic_ops.py b/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_logic_ops.py
index eb5bac45f9..df899124a1 100644
--- a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_logic_ops.py
+++ b/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_logic_ops.py
@@ -9,7 +9,7 @@
 import pytest
 
 from numba_dpex import dpjit
-from numba_dpex.tests._helper import assert_auto_offloading, filter_strings
+from numba_dpex.tests._helper import filter_strings
 
 """ Following cases, dpnp raises NotImplementedError"""
 
diff --git a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_transcedental_functions.py b/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_transcedental_functions.py
index f43d59752f..a2a2d65f46 100644
--- a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_transcedental_functions.py
+++ b/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_transcedental_functions.py
@@ -9,11 +9,7 @@
 import pytest
 
 from numba_dpex import dpjit
-from numba_dpex.tests._helper import (
-    assert_auto_offloading,
-    filter_strings,
-    is_gen12,
-)
+from numba_dpex.tests._helper import filter_strings, is_gen12
 
 """dpnp raise error on : mod, abs and remainder(float32)"""
 list_of_binary_ops = [
diff --git a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_trigonometric_functions.py b/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_trigonometric_functions.py
index 73ef3b4f5d..6c4e877f54 100644
--- a/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_trigonometric_functions.py
+++ b/numba_dpex/tests/dpjit_tests/parfors/test_dpnp_trigonometric_functions.py
@@ -9,11 +9,7 @@
 import pytest
 
 from numba_dpex import dpjit
-from numba_dpex.tests._helper import (
-    assert_auto_offloading,
-    filter_strings,
-    is_gen12,
-)
+from numba_dpex.tests._helper import filter_strings, is_gen12
 
 list_of_filter_strs = [
     "opencl:gpu:0",
diff --git a/numba_dpex/tests/kernel_tests/test_compute_follows_data.py b/numba_dpex/tests/kernel_tests/test_compute_follows_data.py
deleted file mode 100644
index b6e85b7b1d..0000000000
--- a/numba_dpex/tests/kernel_tests/test_compute_follows_data.py
+++ /dev/null
@@ -1,238 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import warnings
-
-import dpctl
-import dpctl.tensor as dpt
-import numpy as np
-import pytest
-
-import numba_dpex
-from numba_dpex.core.exceptions import ComputeFollowsDataInferenceError
-from numba_dpex.tests._helper import (
-    filter_strings,
-    skip_no_level_zero_gpu,
-    skip_no_opencl_gpu,
-)
-
-global_size = 10
-local_size = 1
-N = global_size * local_size
-
-
-@numba_dpex.kernel
-def sum_kernel(a, b, c):
-    i = numba_dpex.get_global_id(0)
-    c[i] = a[i] + b[i]
-
-
-list_of_uniform_types = [
-    (np.array, np.array, np.array),
-    (
-        dpctl.tensor.usm_ndarray,
-        dpctl.tensor.usm_ndarray,
-        dpctl.tensor.usm_ndarray,
-    ),
-]
-
-list_of_dtypes = [
-    np.float32,
-    np.float64,
-]
-
-
-@pytest.fixture(params=list_of_dtypes)
-def input_arrays(request):
-    a = np.array(np.random.random(N), request.param)
-    b = np.array(np.random.random(N), request.param)
-    c = np.zeros_like(a)
-    return a, b, c
-
-
-@pytest.mark.parametrize("offload_device", filter_strings)
-def test_usm_ndarray_argtype(offload_device, input_arrays):
-    usm_type = "device"
-
-    a, b, expected = input_arrays
-    got = np.ones_like(a)
-
-    device = dpctl.SyclDevice(offload_device)
-    queue = dpctl.SyclQueue(device)
-
-    da = dpt.usm_ndarray(
-        a.shape,
-        dtype=a.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue},
-    )
-    da.usm_data.copy_from_host(a.reshape((-1)).view("|u1"))
-
-    db = dpt.usm_ndarray(
-        b.shape,
-        dtype=b.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue},
-    )
-    db.usm_data.copy_from_host(b.reshape((-1)).view("|u1"))
-
-    dc = dpt.usm_ndarray(
-        got.shape,
-        dtype=got.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue},
-    )
-
-    sum_kernel[global_size, local_size](da, db, dc)
-
-    dc.usm_data.copy_to_host(got.reshape((-1)).view("|u1"))
-
-    expected = a + b
-
-    assert np.array_equal(got, expected)
-
-
-@pytest.mark.parametrize("offload_device", filter_strings)
-def test_ndarray_argtype(offload_device, input_arrays):
-    a, b, expected = input_arrays
-    got = np.ones_like(a)
-
-    with numba_dpex.offload_to_sycl_device(offload_device):
-        sum_kernel[global_size, local_size](a, b, got)
-
-    expected = a + b
-
-    assert np.array_equal(got, expected)
-
-
-@pytest.mark.parametrize("offload_device", filter_strings)
-def test_mix_argtype(offload_device, input_arrays):
-    usm_type = "device"
-
-    a, b, _ = input_arrays
-    got = np.ones_like(a)
-
-    device = dpctl.SyclDevice(offload_device)
-    queue = dpctl.SyclQueue(device)
-
-    da = dpt.usm_ndarray(
-        a.shape,
-        dtype=a.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue},
-    )
-    da.usm_data.copy_from_host(a.reshape((-1)).view("|u1"))
-
-    dc = dpt.usm_ndarray(
-        got.shape,
-        dtype=got.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue},
-    )
-
-    with pytest.raises(ComputeFollowsDataInferenceError):
-        sum_kernel[global_size, local_size](da, b, dc)
-
-
-@pytest.mark.parametrize("offload_device", filter_strings)
-def test_context_manager_with_usm_ndarray(offload_device, input_arrays):
-    usm_type = "device"
-
-    a, b, expected = input_arrays
-    got = np.ones_like(a)
-
-    device = dpctl.SyclDevice(offload_device)
-    queue = dpctl.SyclQueue(device)
-
-    da = dpt.usm_ndarray(
-        a.shape,
-        dtype=a.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue},
-    )
-    da.usm_data.copy_from_host(a.reshape((-1)).view("|u1"))
-
-    db = dpt.usm_ndarray(
-        b.shape,
-        dtype=b.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue},
-    )
-    db.usm_data.copy_from_host(b.reshape((-1)).view("|u1"))
-
-    dc = dpt.usm_ndarray(
-        got.shape,
-        dtype=got.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue},
-    )
-
-    with pytest.warns(Warning) as warning:
-        with numba_dpex.offload_to_sycl_device(offload_device):
-            sum_kernel[global_size, local_size](da, db, dc)
-        if not warning:
-            pytest.fail("Warning expected!")
-
-    sum_kernel[global_size, local_size](da, db, dc)
-
-    dc.usm_data.copy_to_host(got.reshape((-1)).view("|u1"))
-
-    expected = a + b
-
-    assert np.array_equal(got, expected)
-
-
-@skip_no_level_zero_gpu
-@skip_no_opencl_gpu
-def test_equivalent_usm_ndarray(input_arrays):
-    usm_type = "device"
-
-    a, b, expected = input_arrays
-    got = np.ones_like(a)
-
-    device1 = dpctl.SyclDevice("level_zero:gpu")
-    queue1 = dpctl.SyclQueue(device1)
-
-    device2 = dpctl.SyclDevice("opencl:gpu")
-    queue2 = dpctl.SyclQueue(device2)
-
-    da = dpt.usm_ndarray(
-        a.shape,
-        dtype=a.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue1},
-    )
-    da.usm_data.copy_from_host(a.reshape((-1)).view("|u1"))
-
-    not_equivalent_db = dpt.usm_ndarray(
-        b.shape,
-        dtype=b.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue2},
-    )
-    not_equivalent_db.usm_data.copy_from_host(b.reshape((-1)).view("|u1"))
-
-    equivalent_db = dpt.usm_ndarray(
-        b.shape,
-        dtype=b.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue1},
-    )
-    equivalent_db.usm_data.copy_from_host(b.reshape((-1)).view("|u1"))
-
-    dc = dpt.usm_ndarray(
-        got.shape,
-        dtype=got.dtype,
-        buffer=usm_type,
-        buffer_ctor_kwargs={"queue": queue1},
-    )
-
-    with pytest.raises(ComputeFollowsDataInferenceError):
-        sum_kernel[global_size, local_size](da, not_equivalent_db, dc)
-
-    sum_kernel[global_size, local_size](da, equivalent_db, dc)
-    dc.usm_data.copy_to_host(got.reshape((-1)).view("|u1"))
-
-    expected = a + b
-    assert np.array_equal(got, expected)
diff --git a/numba_dpex/tests/njit_tests/__init__.py b/numba_dpex/tests/njit_tests/__init__.py
deleted file mode 100644
index 5ca1120d9e..0000000000
--- a/numba_dpex/tests/njit_tests/__init__.py
+++ /dev/null
@@ -1,6 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-from .._helper import *
-from . import *
diff --git a/numba_dpex/tests/njit_tests/dpnp/__init__.py b/numba_dpex/tests/njit_tests/dpnp/__init__.py
deleted file mode 100644
index 141a9eccf2..0000000000
--- a/numba_dpex/tests/njit_tests/dpnp/__init__.py
+++ /dev/null
@@ -1,5 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-from . import *
diff --git a/numba_dpex/tests/njit_tests/dpnp/_helper.py b/numba_dpex/tests/njit_tests/dpnp/_helper.py
deleted file mode 100644
index b8d3910719..0000000000
--- a/numba_dpex/tests/njit_tests/dpnp/_helper.py
+++ /dev/null
@@ -1,23 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import string
-
-
-def wrapper_function(args, code, globals, function_name="func"):
-    function_text = f"""\
-def {function_name}({args}):
-    return {code}
-"""
-    return compile_function(function_text, function_name, globals)
-
-
-def args_string(args_count):
-    return ", ".join(list(string.ascii_lowercase[:args_count]))
-
-
-def compile_function(function_text, function_name, globals):
-    locals = {}
-    exec(function_text, globals, locals)
-    return locals[function_name]
diff --git a/numba_dpex/tests/njit_tests/dpnp/test_numpy_array_ops.py b/numba_dpex/tests/njit_tests/dpnp/test_numpy_array_ops.py
deleted file mode 100644
index 064a61f0b6..0000000000
--- a/numba_dpex/tests/njit_tests/dpnp/test_numpy_array_ops.py
+++ /dev/null
@@ -1,142 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import numpy as np
-import pytest
-from numba import njit
-
-from numba_dpex.tests._helper import (
-    dpnp_debug,
-    filter_strings,
-    is_gen12,
-    skip_no_dpnp,
-    skip_windows,
-)
-
-from ._helper import wrapper_function
-
-pytestmark = skip_no_dpnp
-
-list_of_dtypes = [
-    np.int32,
-    np.int64,
-    np.float32,
-    np.float64,
-]
-
-
-@pytest.fixture(params=list_of_dtypes)
-def input_arrays(request):
-    # The size of input and out arrays to be used
-    N = 10
-    a = np.array(np.random.random(N), request.param)
-    b = np.array(np.random.random(N), request.param)
-    return a, b
-
-
-list_of_shape = [
-    (10),
-    (5, 2),
-]
-
-
-@pytest.fixture(params=list_of_shape)
-def get_shape(request):
-    return request.param
-
-
-list_of_unary_ops = [
-    "sum",
-    "prod",
-    "max",
-    "min",
-    "mean",
-    "argmax",
-    "argmin",
-    "argsort",
-    "copy",
-    "cumsum",
-    "cumprod",
-]
-
-
-@pytest.fixture(params=list_of_unary_ops)
-def unary_op(request):
-    return (
-        wrapper_function("a", f"a.{request.param}()", globals()),
-        request.param,
-    )
-
-
-@pytest.mark.parametrize("filter_str", filter_strings)
-def test_unary_ops(filter_str, unary_op, input_arrays, get_shape, capfd):
-    a = input_arrays[0]
-    op, name = unary_op
-    if name != "argsort" and name != "copy":
-        a = np.reshape(a, get_shape)
-    if name == "cumprod" and (
-        filter_str == "opencl:cpu:0"
-        or a.dtype == np.int32
-        or is_gen12(filter_str)
-    ):
-        pytest.skip()
-    if name == "cumsum" and (
-        filter_str == "opencl:cpu:0"
-        or a.dtype == np.int32
-        or is_gen12(filter_str)
-    ):
-        pytest.skip()
-    if name == "mean" and is_gen12(filter_str):
-        pytest.skip()
-    if name == "argmax" and is_gen12(filter_str):
-        pytest.skip()
-
-    actual = np.empty(shape=a.shape, dtype=a.dtype)
-    expected = np.empty(shape=a.shape, dtype=a.dtype)
-
-    f = njit(op)
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(a)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-    expected = op(a)
-    np.testing.assert_allclose(actual, expected, rtol=1e-3, atol=0)
-
-
-list_of_indices = [
-    np.array([0, 2, 5], dtype=np.int64),
-    np.array([0, 5], dtype=np.int32),
-]
-
-
-@pytest.fixture(params=list_of_indices)
-def indices(request):
-    return request.param
-
-
-def get_take_fn():
-    return wrapper_function("a, ind", "a.take(ind)", globals())
-
-
-@skip_windows
-@pytest.mark.parametrize("filter_str", filter_strings)
-def test_take(filter_str, input_arrays, indices, capfd):
-    a = input_arrays[0]
-    fn = get_take_fn()
-
-    actual = np.empty(shape=a.shape, dtype=a.dtype)
-    expected = np.empty(shape=a.shape, dtype=a.dtype)
-
-    f = njit(fn)
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(a, indices)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-    expected = fn(a, indices)
-    np.testing.assert_allclose(actual, expected, rtol=1e-3, atol=0)
diff --git a/numba_dpex/tests/njit_tests/dpnp/test_numpy_indexing.py b/numba_dpex/tests/njit_tests/dpnp/test_numpy_indexing.py
deleted file mode 100644
index 955d6268d2..0000000000
--- a/numba_dpex/tests/njit_tests/dpnp/test_numpy_indexing.py
+++ /dev/null
@@ -1,63 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import numpy as np
-import pytest
-from numba import njit
-
-from numba_dpex.tests._helper import dpnp_debug, filter_strings, skip_no_dpnp
-
-pytestmark = skip_no_dpnp
-
-
-@pytest.mark.parametrize("filter_str", filter_strings)
-@pytest.mark.parametrize("offset", [0, 1], ids=["0", "1"])
-@pytest.mark.parametrize(
-    "array",
-    [
-        [[0, 0], [0, 0]],
-        [[1, 2], [1, 2]],
-        [[1, 2], [3, 4]],
-        [[0, 1, 2], [3, 4, 5], [6, 7, 8]],
-        [[0, 1, 2, 3, 4], [5, 6, 7, 8, 9]],
-        [[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]],
-        [
-            [[[1, 2], [3, 4]], [[1, 2], [2, 1]]],
-            [[[1, 3], [3, 1]], [[0, 1], [1, 3]]],
-        ],
-        [
-            [[[1, 2, 3], [3, 4, 5]], [[1, 2, 3], [2, 1, 0]]],
-            [[[1, 3, 5], [3, 1, 0]], [[0, 1, 2], [1, 3, 4]]],
-        ],
-        [
-            [[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]],
-            [[[13, 14, 15], [16, 17, 18]], [[19, 20, 21], [22, 23, 24]]],
-        ],
-    ],
-    ids=[
-        "[[0, 0], [0, 0]]",
-        "[[1, 2], [1, 2]]",
-        "[[1, 2], [3, 4]]",
-        "[[0, 1, 2], [3, 4, 5], [6, 7, 8]]",
-        "[[0, 1, 2, 3, 4], [5, 6, 7, 8, 9]]",
-        "[[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]]",
-        "[[[[1, 2], [3, 4]], [[1, 2], [2, 1]]], [[[1, 3], [3, 1]], [[0, 1], [1, 3]]]]",
-        "[[[[1, 2, 3], [3, 4, 5]], [[1, 2, 3], [2, 1, 0]]], [[[1, 3, 5], [3, 1, 0]], [[0, 1, 2], [1, 3, 4]]]]",
-        "[[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]], [[[13, 14, 15], [16, 17, 18]], [[19, 20, 21], [22, 23, 24]]]]",
-    ],
-)
-def test_diagonal(array, offset, filter_str):
-    a = np.array(array)
-
-    def fn(a, offset):
-        return np.diagonal(a, offset)
-
-    f = njit(fn)
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(a, offset)
-
-    expected = fn(a, offset)
-    np.testing.assert_allclose(actual, expected, rtol=1e-3, atol=0)
diff --git a/numba_dpex/tests/njit_tests/dpnp/test_numpy_linalg.py b/numba_dpex/tests/njit_tests/dpnp/test_numpy_linalg.py
deleted file mode 100644
index ccb99ea638..0000000000
--- a/numba_dpex/tests/njit_tests/dpnp/test_numpy_linalg.py
+++ /dev/null
@@ -1,326 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import numpy as np
-import pytest
-from numba import njit
-
-from numba_dpex.tests._helper import (
-    dpnp_debug,
-    filter_strings,
-    filter_strings_with_skips_for_opencl,
-    skip_no_dpnp,
-)
-
-from ._helper import args_string, wrapper_function
-
-pytestmark = skip_no_dpnp
-
-
-# From https://github.com/IntelPython/dpnp/blob/0.4.0/tests/test_linalg.py#L8
-def vvsort(val, vec):
-    size = val.size
-    for i in range(size):
-        imax = i
-        for j in range(i + 1, size):
-            if np.abs(val[imax]) < np.abs(val[j]):
-                imax = j
-
-        temp = val[i]
-        val[i] = val[imax]
-        val[imax] = temp
-
-        if not (vec is None):
-            for k in range(size):
-                temp = vec[k, i]
-                vec[k, i] = vec[k, imax]
-                vec[k, imax] = temp
-
-
-def get_fn(name, nargs):
-    args = args_string(nargs)
-    return wrapper_function(args, f"np.{name}({args})", globals())
-
-
-list_of_dtypes = [
-    np.int32,
-    np.int64,
-    np.float32,
-    np.float64,
-]
-
-
-@pytest.fixture(params=list_of_dtypes)
-def eig_input(request):
-    # The size of input and out arrays to be used
-    N = 10
-    a = np.arange(N * N, dtype=request.param).reshape((N, N))
-    symm_a = (
-        np.tril(a)
-        + np.tril(a, -1).T
-        + +np.diag(np.full((N,), N * N, dtype=request.param))
-    )
-    return symm_a
-
-
-@pytest.mark.parametrize("filter_str", filter_strings_with_skips_for_opencl)
-def test_eig(filter_str, eig_input, capfd):
-    a = eig_input
-    fn = get_fn("linalg.eig", 1)
-    f = njit(fn)
-
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual_val, actual_vec = f(a)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-        expected_val, expected_vec = fn(a)
-
-        # sort val/vec by abs value
-        vvsort(actual_val, actual_vec)
-        vvsort(expected_val, expected_vec)
-
-        # NP change sign of vectors
-        for i in range(expected_vec.shape[1]):
-            if expected_vec[0, i] * actual_vec[0, i] < 0:
-                expected_vec[:, i] = -expected_vec[:, i]
-
-        assert np.allclose(actual_val, expected_val)
-        assert np.allclose(actual_vec, expected_vec)
-
-
-@pytest.fixture(params=list_of_dtypes)
-def dtype(request):
-    return request.param
-
-
-list_of_dim = [
-    (10, 1, 10, 1),
-    (10, 1, 10, 2),
-    (2, 10, 10, 1),
-    (10, 2, 2, 10),
-]
-
-
-@pytest.fixture(params=list_of_dim)
-def dot_input(request):
-    # The size of input and out arrays to be used
-    a1, a2, b1, b2 = request.param
-    a = np.array(np.random.random(a1 * a2))
-    b = np.array(np.random.random(b1 * b2))
-    if a2 != 1:
-        a = a.reshape(a1, a2)
-    if b2 != 1:
-        b = b.reshape(b1, b2)
-
-    return a, b
-
-
-list_of_dot_name = ["dot", "vdot"]
-
-
-@pytest.fixture(params=list_of_dot_name)
-def dot_name(request):
-    return request.param
-
-
-@pytest.mark.parametrize("filter_str", filter_strings_with_skips_for_opencl)
-def test_dot(filter_str, dot_name, dot_input, dtype, capfd):
-    a, b = dot_input
-
-    if dot_name == "vdot":
-        if a.size != b.size:
-            pytest.skip("vdot only supports same sized arrays")
-
-    a = a.astype(dtype)
-    b = b.astype(dtype)
-    fn = get_fn(dot_name, 2)
-    f = njit(fn)
-
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(a, b)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-        expected = fn(a, b)
-        assert np.allclose(actual, expected)
-
-
-@pytest.mark.parametrize("filter_str", filter_strings_with_skips_for_opencl)
-def test_matmul(filter_str, dtype, capfd):
-    a = np.array(np.random.random(10 * 2), dtype=dtype).reshape(10, 2)
-    b = np.array(np.random.random(2 * 10), dtype=dtype).reshape(2, 10)
-    fn = get_fn("matmul", 2)
-    f = njit(fn)
-
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(a, b)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-        expected = fn(a, b)
-        assert np.allclose(actual, expected)
-
-
-@pytest.mark.parametrize("filter_str", filter_strings)
-@pytest.mark.skip(reason="dpnp does not support it yet")
-def test_cholesky(filter_str, dtype, capfd):
-    a = np.array([[1, -2], [2, 5]], dtype=dtype)
-    fn = get_fn("linalg.cholesky", 1)
-    f = njit(fn)
-
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(a)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-        expected = fn(a)
-        assert np.allclose(actual, expected)
-
-
-list_of_det_input = [
-    [[0, 0], [0, 0]],
-    [[1, 2], [1, 2]],
-    [[1, 2], [3, 4]],
-    [[[1, 2], [3, 4]], [[1, 2], [2, 1]], [[1, 3], [3, 1]]],
-    [
-        [[[1, 2], [3, 4]], [[1, 2], [2, 1]]],
-        [[[1, 3], [3, 1]], [[0, 1], [1, 3]]],
-    ],
-]
-
-
-@pytest.fixture(params=list_of_det_input)
-def det_input(request):
-    return request.param
-
-
-@pytest.mark.parametrize("filter_str", filter_strings)
-def test_det(filter_str, det_input, dtype, capfd):
-    a = np.array(det_input, dtype=dtype)
-    fn = get_fn("linalg.det", 1)
-    f = njit(fn)
-
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(a)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-        expected = fn(a)
-        assert np.allclose(actual, expected)
-
-
-@pytest.mark.parametrize("filter_str", filter_strings_with_skips_for_opencl)
-def test_multi_dot(filter_str, capfd):
-    def fn(A, B, C, D):
-        c = np.linalg.multi_dot([A, B, C, D])
-        return c
-
-    A = np.random.random((10000, 100))
-    B = np.random.random((100, 1000))
-    C = np.random.random((1000, 5))
-    D = np.random.random((5, 333))
-    f = njit(fn)
-
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(A, B, C, D)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-        expected = fn(A, B, C, D)
-        assert np.allclose(actual, expected)
-
-
-list_of_power = [2, 3, 0]
-
-
-@pytest.fixture(params=list_of_power)
-def power(request):
-    return request.param
-
-
-list_of_matrix_power_input = [
-    [[0, 0], [0, 0]],
-    [[1, 2], [1, 2]],
-    [[1, 2], [3, 4]],
-]
-
-
-@pytest.fixture(params=list_of_matrix_power_input)
-def matrix_power_input(request):
-    return request.param
-
-
-@pytest.mark.parametrize("filter_str", filter_strings_with_skips_for_opencl)
-def test_matrix_power(filter_str, matrix_power_input, power, dtype, capfd):
-    a = np.array(matrix_power_input, dtype=dtype)
-    fn = get_fn("linalg.matrix_power", 2)
-    f = njit(fn)
-
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(a, power)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-        expected = fn(a, power)
-        assert np.allclose(actual, expected)
-
-
-@pytest.mark.parametrize("filter_str", filter_strings)
-@pytest.mark.parametrize(
-    "matrix_rank_input",
-    [
-        pytest.param(
-            np.eye(4),
-            marks=pytest.mark.xfail(reason="dpnp does not support it yet"),
-        ),
-        np.ones((4,)),
-        pytest.param(
-            np.ones((4, 4)),
-            marks=pytest.mark.xfail(reason="dpnp does not support it yet"),
-        ),
-        np.zeros((4,)),
-    ],
-)
-def test_matrix_rank(filter_str, matrix_rank_input, capfd):
-    fn = get_fn("linalg.matrix_rank", 1)
-    f = njit(fn)
-
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(matrix_rank_input)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-        expected = fn(matrix_rank_input)
-        assert np.allclose(actual, expected)
-
-
-@pytest.mark.parametrize("filter_str", filter_strings_with_skips_for_opencl)
-def test_eigvals(filter_str, eig_input, capfd):
-    a = eig_input
-    fn = get_fn("linalg.eigvals", 1)
-    f = njit(fn)
-
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual_val = f(a)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-        expected_val = fn(a)
-
-        # sort val/vec by abs value
-        vvsort(actual_val, None)
-        vvsort(expected_val, None)
-
-        assert np.allclose(actual_val, expected_val)
diff --git a/numba_dpex/tests/njit_tests/dpnp/test_numpy_logic.py b/numba_dpex/tests/njit_tests/dpnp/test_numpy_logic.py
deleted file mode 100644
index e91860adcc..0000000000
--- a/numba_dpex/tests/njit_tests/dpnp/test_numpy_logic.py
+++ /dev/null
@@ -1,48 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import numpy as np
-import pytest
-from numba import njit
-
-from numba_dpex.tests._helper import dpnp_debug, filter_strings, skip_no_dpnp
-
-pytestmark = skip_no_dpnp
-
-
-@pytest.mark.parametrize("filter_str", filter_strings)
-@pytest.mark.parametrize(
-    "dtype", [np.bool_, np.int32, np.int64, np.float32, np.float64]
-)
-@pytest.mark.parametrize(
-    "shape",
-    [(0,), (4,), (2, 3)],
-    ids=["(0,)", "(4,)", "(2, 3)"],
-)
-def test_all(dtype, shape, filter_str):
-    size = 1
-    for i in range(len(shape)):
-        size *= shape[i]
-
-    for i in range(2**size):
-        t = i
-
-        a = np.empty(size, dtype=dtype)
-
-        for j in range(size):
-            a[j] = 0 if t % 2 == 0 else j + 1
-            t = t >> 1
-
-        a = a.reshape(shape)
-
-        def fn(a):
-            return np.all(a)
-
-        f = njit(fn)
-        with dpctl.device_context(filter_str), dpnp_debug():
-            actual = f(a)
-
-        expected = fn(a)
-        np.testing.assert_allclose(actual, expected, rtol=1e-3, atol=0)
diff --git a/numba_dpex/tests/njit_tests/dpnp/test_numpy_manipulation.py b/numba_dpex/tests/njit_tests/dpnp/test_numpy_manipulation.py
deleted file mode 100644
index 00d29e16a9..0000000000
--- a/numba_dpex/tests/njit_tests/dpnp/test_numpy_manipulation.py
+++ /dev/null
@@ -1,33 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import numpy as np
-import pytest
-from numba import njit
-
-from numba_dpex.tests._helper import dpnp_debug, filter_strings, skip_no_dpnp
-
-pytestmark = skip_no_dpnp
-
-
-@pytest.mark.parametrize("filter_str", filter_strings)
-@pytest.mark.parametrize(
-    "arr",
-    [[1, 2, 3, 4], [1, 2, 3, 4, 5, 6, 7, 8, 9]],
-    ids=["[1, 2, 3, 4]", "[1, 2, 3, 4, 5, 6, 7, 8, 9]"],
-)
-def test_repeat(filter_str, arr):
-    a = np.array(arr)
-    repeats = 2
-
-    def fn(a, repeats):
-        return np.repeat(a, repeats)
-
-    f = njit(fn)
-    with dpctl.device_context(filter_str), dpnp_debug():
-        actual = f(a, repeats)
-
-    expected = fn(a, repeats)
-    np.testing.assert_allclose(actual, expected, rtol=1e-3, atol=0)
diff --git a/numba_dpex/tests/njit_tests/dpnp/test_numpy_rng.py b/numba_dpex/tests/njit_tests/dpnp/test_numpy_rng.py
deleted file mode 100644
index 0bb2a3e1ff..0000000000
--- a/numba_dpex/tests/njit_tests/dpnp/test_numpy_rng.py
+++ /dev/null
@@ -1,229 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import numpy as np
-import pytest
-from numba import njit
-
-from numba_dpex.tests._helper import (
-    dpnp_debug,
-    filter_strings_level_zero_gpu,
-    skip_no_dpnp,
-)
-
-from ._helper import wrapper_function
-
-pytestmark = skip_no_dpnp
-
-# dpnp throws -30 (CL_INVALID_VALUE) when invoked with multiple kinds of
-# devices at runtime, so testing for level_zero only
-
-
-list_of_size = [
-    9,
-    (2, 5),
-    (3, 2, 4),
-]
-
-none_size = [None]
-
-
-@pytest.fixture(params=list_of_size)
-def unary_size(request):
-    return request.param
-
-
-@pytest.fixture(params=list_of_size + none_size)
-def three_arg_size(request):
-    return request.param
-
-
-list_of_one_arg = [
-    ("random_sample", 0.0, 1.0),
-    ("ranf", 0.0, 1.0),
-    ("sample", 0.0, 1.0),
-    ("random", 0.0, 1.0),
-    ("standard_exponential", 0.0, None),
-    ("standard_normal", None, None),
-    ("standard_cauchy", None, None),
-]
-
-
-@pytest.fixture(params=list_of_one_arg)
-def one_arg_fn(request):
-    function = wrapper_function(
-        "size", f"np.random.{request.param[0]}(size)", globals()
-    )
-    return function, request.param
-
-
-@pytest.mark.parametrize("filter_str", filter_strings_level_zero_gpu)
-def test_one_arg_fn(filter_str, one_arg_fn, unary_size, capfd):
-    op, params = one_arg_fn
-    name, low, high = params
-    f = njit(op)
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(unary_size)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-        if low is not None:
-            assert np.all(actual >= low)
-        if high is not None:
-            assert np.all(actual < high)
-
-
-list_of_two_arg_fn = [
-    ("chisquare", 3, 0, None),
-    ("exponential", 3.0, 0, None),
-    ("gamma", 2.0, 0, None),
-    ("geometric", 0.35, 0, None),
-    ("poisson", 5.0, 0, None),
-    ("rayleigh", 2.0, 0, None),
-    ("standard_gamma", 2.0, 0, None),
-    ("weibull", 5.0, 0, None),
-]
-
-
-@pytest.fixture(params=list_of_two_arg_fn)
-def two_arg_fn(request):
-    return request.param
-
-
-def get_two_arg_fn(op_name):
-    return wrapper_function("a, b", f"np.random.{op_name}(a, b)", globals())
-
-
-@pytest.mark.parametrize("filter_str", filter_strings_level_zero_gpu)
-def test_two_arg_fn(filter_str, two_arg_fn, unary_size, capfd):
-    op_name, first_arg, low, high = two_arg_fn
-
-    if op_name == "gamma":
-        pytest.skip(
-            "AttributeError: 'NoneType' object has no attribute 'ravel'"
-        )
-    op = get_two_arg_fn(op_name)
-    f = njit(op)
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(first_arg, unary_size)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-        if low is not None and high is None:
-            if np.isscalar(actual):
-                assert actual >= low
-            else:
-                actual = actual.ravel()
-                assert np.all(actual >= low)
-
-
-list_of_three_arg_fn = [
-    ("randint", 2, 23, 0, None),
-    ("random_integers", 2, 23, 1, None),
-    ("beta", 2.56, 0.8, 0, 1.0),
-    ("binomial", 5, 0.0, 0, 1.0),
-    ("gumbel", 0.5, 0.1, None, None),
-    ("laplace", 0.0, 1.0, None, None),
-    ("lognormal", 3.0, 1.0, None, None),
-    ("multinomial", 100, np.array([1 / 7.0] * 5), 0, 100),
-    ("multivariate_normal", (1, 2), [[1, 0], [0, 1]], None, None),
-    ("negative_binomial", 1, 0.1, 0, None),
-    ("normal", 0.0, 0.1, None, None),
-    ("uniform", -1.0, 0.0, -1.0, 0.0),
-]
-
-
-@pytest.fixture(params=list_of_three_arg_fn)
-def three_arg_fn(request):
-    return request.param
-
-
-def get_three_arg_fn(op_name):
-    return wrapper_function(
-        "a, b, c", f"np.random.{op_name}(a, b, c)", globals()
-    )
-
-
-@pytest.mark.parametrize("filter_str", filter_strings_level_zero_gpu)
-def test_three_arg_fn(filter_str, three_arg_fn, three_arg_size, capfd):
-    op_name, first_arg, second_arg, low, high = three_arg_fn
-
-    if op_name == "multinomial":
-        pytest.skip("DPNP RNG Error: dpnp_rng_multinomial_c() failed")
-    elif op_name == "multivariate_normal":
-        pytest.skip(
-            "No implementation of function Function(<class "
-            "'numba_dpex.dpnp_iface.stubs.dpnp.multivariate_normal'>) found for signature"
-        )
-    elif op_name == "negative_binomial":
-        pytest.skip("DPNP RNG Error: dpnp_rng_negative_binomial_c() failed.")
-    elif op_name == "gumbel":
-        pytest.skip("DPNP error")
-
-    op = get_three_arg_fn(op_name)
-    f = njit(op)
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(first_arg, second_arg, three_arg_size)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-        if low is not None and high is None:
-            if second_arg:
-                low = first_arg
-                high = second_arg
-                assert np.all(actual >= low)
-                assert np.all(actual <= high)
-            else:
-                high = first_arg
-                assert np.all(actual >= low)
-                assert np.all(actual <= high)
-        elif low is not None and high is not None:
-            if np.isscalar(actual):
-                assert actual >= low
-                assert actual <= high
-            else:
-                actual = actual.ravel()
-                assert np.all(actual >= low)
-                assert np.all(actual <= high)
-
-
-@pytest.mark.parametrize("filter_str", filter_strings_level_zero_gpu)
-def test_rand(filter_str):
-    @njit
-    def f():
-        c = np.random.rand(3, 2)
-        return c
-
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f()
-
-        actual = actual.ravel()
-        assert np.all(actual >= 0.0)
-        assert np.all(actual < 1.0)
-
-
-@pytest.mark.parametrize("filter_str", filter_strings_level_zero_gpu)
-def test_hypergeometric(filter_str, three_arg_size):
-    @njit
-    def f(ngood, nbad, nsamp, size):
-        res = np.random.hypergeometric(ngood, nbad, nsamp, size)
-        return res
-
-    ngood, nbad, nsamp = 100, 2, 10
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(ngood, nbad, nsamp, three_arg_size)
-
-        if np.isscalar(actual):
-            assert actual >= 0
-            assert actual <= min(nsamp, ngood + nbad)
-        else:
-            actual = actual.ravel()
-            assert np.all(actual >= 0)
-            assert np.all(actual <= min(nsamp, ngood + nbad))
diff --git a/numba_dpex/tests/njit_tests/dpnp/test_numpy_sort_search_count.py b/numba_dpex/tests/njit_tests/dpnp/test_numpy_sort_search_count.py
deleted file mode 100644
index 13862b8bb0..0000000000
--- a/numba_dpex/tests/njit_tests/dpnp/test_numpy_sort_search_count.py
+++ /dev/null
@@ -1,117 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import numpy as np
-import pytest
-from numba import njit
-
-from numba_dpex.tests._helper import dpnp_debug, filter_strings, skip_no_dpnp
-
-from ._helper import wrapper_function
-
-pytestmark = skip_no_dpnp
-
-list_of_dtypes = [
-    np.int32,
-    np.int64,
-    np.float32,
-    np.float64,
-]
-
-
-@pytest.fixture(params=list_of_dtypes)
-def input_arrays(request):
-    # The size of input and out arrays to be used
-    N = 10
-    a = np.array(np.random.random(N), request.param)
-    b = np.array(np.random.random(N), request.param)
-    return a, b
-
-
-list_of_shape = [
-    (10),
-    (5, 2),
-]
-
-
-@pytest.fixture(params=list_of_shape)
-def get_shape(request):
-    return request.param
-
-
-list_of_unary_ops = [
-    "sort",
-    "argmax",
-    "argmin",
-    "argsort",
-]
-
-
-@pytest.fixture(params=list_of_unary_ops)
-def unary_op(request):
-    return (
-        wrapper_function("a", f"np.{request.param}(a)", globals()),
-        request.param,
-    )
-
-
-@pytest.mark.parametrize("filter_str", filter_strings)
-def test_unary_ops(filter_str, unary_op, input_arrays, get_shape, capfd):
-    a = input_arrays[0]
-    op, name = unary_op
-    if name != "argsort" and name != "sort":
-        a = np.reshape(a, get_shape)
-    actual = np.empty(shape=a.shape, dtype=a.dtype)
-    expected = np.empty(shape=a.shape, dtype=a.dtype)
-
-    f = njit(op)
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(a)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-    expected = op(a)
-    np.testing.assert_allclose(actual, expected, rtol=1e-3, atol=0)
-
-
-@pytest.mark.skip(reason="Flaky tests")
-@pytest.mark.parametrize("filter_str", filter_strings)
-@pytest.mark.parametrize("kth", [0, 1], ids=["0", "1"])
-@pytest.mark.parametrize(
-    "array",
-    [
-        [3, 4, 2, 1],
-        [[1, 0], [3, 0]],
-        [[3, 2], [1, 6]],
-        [[4, 2, 3], [3, 4, 1]],
-        [[[1, -3], [3, 0]], [[5, 2], [0, 1]], [[1, 0], [0, 1]]],
-        [
-            [[[8, 2], [3, 0]], [[5, 2], [0, 1]]],
-            [[[1, 3], [3, 1]], [[5, 2], [0, 1]]],
-        ],
-    ],
-    ids=[
-        "[3, 4, 2, 1]",
-        "[[1, 0], [3, 0]]",
-        "[[3, 2], [1, 6]]",
-        "[[4, 2, 3], [3, 4, 1]]",
-        "[[[1, -3], [3, 0]], [[5, 2], [0, 1]], [[1, 0], [0, 1]]]",
-        "[[[[8, 2], [3, 0]], [[5, 2], [0, 1]]], [[[1, 3], [3, 1]], [[5, 2], [0, 1]]]]",
-    ],
-)
-def test_partition(array, kth, filter_str):
-    a = np.array(array)
-
-    def fn(a, kth):
-        return np.partition(a, kth)
-
-    f = njit(fn)
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(a, kth)
-
-    expected = fn(a, kth)
-    np.testing.assert_allclose(actual, expected, rtol=1e-3, atol=0)
diff --git a/numba_dpex/tests/njit_tests/dpnp/test_numpy_statistics.py b/numba_dpex/tests/njit_tests/dpnp/test_numpy_statistics.py
deleted file mode 100644
index 77bdc0b00a..0000000000
--- a/numba_dpex/tests/njit_tests/dpnp/test_numpy_statistics.py
+++ /dev/null
@@ -1,86 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import numpy as np
-import pytest
-from numba import njit
-
-from numba_dpex.tests._helper import (
-    dpnp_debug,
-    filter_strings_with_skips_for_opencl,
-    skip_no_dpnp,
-)
-
-from ._helper import wrapper_function
-
-pytestmark = skip_no_dpnp
-
-list_of_dtypes = [
-    np.int32,
-    np.int64,
-    np.float32,
-    np.float64,
-]
-
-
-@pytest.fixture(params=list_of_dtypes)
-def input_arrays(request):
-    # The size of input and out arrays to be used
-    N = 100
-    a = np.array(np.random.random(N), request.param)
-    b = np.array(np.random.random(N), request.param)
-    return a, b
-
-
-list_of_shape = [
-    (100),
-    (50, 2),
-    (10, 5, 2),
-]
-
-
-@pytest.fixture(params=list_of_shape)
-def get_shape(request):
-    return request.param
-
-
-list_of_unary_ops = [
-    "max",
-    "amax",
-    "min",
-    "amin",
-    "median",
-    "mean",
-    "cov",
-]
-
-
-@pytest.fixture(params=list_of_unary_ops)
-def unary_op(request):
-    return (
-        wrapper_function("a", f"np.{request.param}(a)", globals()),
-        request.param,
-    )
-
-
-@pytest.mark.parametrize("filter_str", filter_strings_with_skips_for_opencl)
-def test_unary_ops(filter_str, unary_op, input_arrays, get_shape, capfd):
-    a = input_arrays[0]
-    op, name = unary_op
-    if name != "cov":
-        a = np.reshape(a, get_shape)
-
-    actual = np.empty(shape=a.shape, dtype=a.dtype)
-    expected = np.empty(shape=a.shape, dtype=a.dtype)
-
-    f = njit(op)
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(a)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-    expected = op(a)
-    np.testing.assert_allclose(actual, expected, rtol=1e-3, atol=0)
diff --git a/numba_dpex/tests/njit_tests/dpnp/test_numpy_transcendentals.py b/numba_dpex/tests/njit_tests/dpnp/test_numpy_transcendentals.py
deleted file mode 100644
index cba8e058ed..0000000000
--- a/numba_dpex/tests/njit_tests/dpnp/test_numpy_transcendentals.py
+++ /dev/null
@@ -1,138 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import numpy as np
-import pytest
-from numba import njit
-
-from numba_dpex.tests._helper import (
-    dpnp_debug,
-    filter_strings,
-    is_gen12,
-    skip_no_dpnp,
-)
-
-from ._helper import wrapper_function
-
-pytestmark = skip_no_dpnp
-
-list_of_int_dtypes = [
-    np.int32,
-    np.int64,
-]
-
-
-list_of_float_dtypes = [
-    np.float32,
-    np.float64,
-]
-
-
-@pytest.fixture(params=list_of_int_dtypes + list_of_float_dtypes)
-def input_array(request):
-    # The size of input and out arrays to be used
-    N = 100
-    a = np.array(np.random.random(N), request.param)
-    return a
-
-
-@pytest.fixture(params=list_of_float_dtypes)
-def input_nan_array(request):
-    # The size of input and out arrays to be used
-    N = 100
-    a = np.array(np.random.random(N), request.param)
-    for i in range(5):
-        a[N - 1 - i] = np.nan
-    return a
-
-
-list_of_shape = [
-    (100),
-    (50, 2),
-    (10, 5, 2),
-]
-
-
-@pytest.fixture(params=list_of_shape)
-def get_shape(request):
-    return request.param
-
-
-list_of_unary_ops = [
-    "sum",
-    "prod",
-    "cumsum",
-    "cumprod",
-]
-
-list_of_nan_ops = [
-    "nansum",
-    "nanprod",
-]
-
-
-def get_func(name):
-    return wrapper_function("a", f"np.{name}(a)", globals())
-
-
-@pytest.fixture(params=list_of_unary_ops + list_of_nan_ops)
-def unary_op(request):
-    fn = get_func(request.param)
-    return fn, request.param
-
-
-@pytest.fixture(params=list_of_nan_ops)
-def unary_nan_op(request):
-    fn = get_func(request.param)
-    return fn, request.param
-
-
-@pytest.mark.parametrize("filter_str", filter_strings)
-def test_unary_ops(filter_str, unary_op, input_array, get_shape, capfd):
-    a = input_array
-    a = np.reshape(a, get_shape)
-    op, name = unary_op
-    if (name == "cumprod" or name == "cumsum") and (
-        filter_str == "opencl:cpu:0" or is_gen12(filter_str)
-    ):
-        pytest.skip()
-    actual = np.empty(shape=a.shape, dtype=a.dtype)
-    expected = np.empty(shape=a.shape, dtype=a.dtype)
-
-    f = njit(op)
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(a)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-    expected = op(a)
-    max_abs_err = np.sum(actual - expected)
-    assert max_abs_err < 1e-4
-
-
-@pytest.mark.parametrize("filter_str", filter_strings)
-def test_unary_nan_ops(
-    filter_str, unary_nan_op, input_nan_array, get_shape, capfd
-):
-    a = input_nan_array
-    a = np.reshape(a, get_shape)
-    op, name = unary_nan_op
-    actual = np.empty(shape=a.shape, dtype=a.dtype)
-    expected = np.empty(shape=a.shape, dtype=a.dtype)
-
-    if name == "nansum" and is_gen12(filter_str):
-        pytest.skip()
-
-    f = njit(op)
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), dpnp_debug():
-        actual = f(a)
-        captured = capfd.readouterr()
-        assert "dpnp implementation" in captured.out
-
-    expected = op(a)
-    max_abs_err = np.sum(actual - expected)
-    assert max_abs_err < 1e-4
diff --git a/numba_dpex/tests/test_controllable_fallback.py b/numba_dpex/tests/test_controllable_fallback.py
deleted file mode 100644
index bdd7e7d983..0000000000
--- a/numba_dpex/tests/test_controllable_fallback.py
+++ /dev/null
@@ -1,75 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import warnings
-
-import dpctl
-import numba
-import numpy as np
-import pytest
-
-from numba_dpex import config
-from numba_dpex.tests._helper import skip_no_opencl_gpu
-
-
-@skip_no_opencl_gpu
-class TestParforFallback:
-    def test_parfor_fallback_true(self):
-        @numba.jit
-        def fill_value(i):
-            return i
-
-        def inner_call_fallback():
-            x = 10
-            a = np.empty(shape=x, dtype=np.float32)
-
-            for i in numba.prange(x):
-                a[i] = fill_value(i)
-
-            return a
-
-        config.DEBUG = 1
-        with warnings.catch_warnings(record=True) as w:
-            device = dpctl.SyclDevice("opencl:gpu")
-            with dpctl.device_context(device):
-                fn = numba.njit(parallel=True)(inner_call_fallback)
-                fallback_true = fn()
-
-        ref_result = inner_call_fallback()
-        config.DEBUG = 0
-
-        np.testing.assert_array_equal(fallback_true, ref_result)
-        assert "Failed to offload parfor" in str(w[-1].message)
-
-    @pytest.mark.xfail
-    def test_parfor_fallback_false(self):
-        @numba.jit
-        def fill_value(i):
-            return i
-
-        def inner_call_fallback():
-            x = 10
-            a = np.empty(shape=x, dtype=np.float32)
-
-            for i in numba.prange(x):
-                a[i] = fill_value(i)
-
-            return a
-
-        try:
-            config.DEBUG = 1
-            config.FALLBACK_ON_CPU = 0
-            with warnings.catch_warnings(record=True) as w:
-                device = dpctl.SyclDevice("opencl:gpu")
-                with dpctl.device_context(device):
-                    fn = numba.njit(parallel=True)(inner_call_fallback)
-                    fallback_false = fn()
-
-        finally:
-            ref_result = inner_call_fallback()
-            config.FALLBACK_ON_CPU = 1
-            config.DEBUG = 0
-
-            not np.testing.assert_array_equal(fallback_false, ref_result)
-            assert "Failed to offload parfor" not in str(w[-1].message)
diff --git a/numba_dpex/tests/test_dpnp_functions.py b/numba_dpex/tests/test_dpnp_functions.py
deleted file mode 100644
index 7ee00b6b81..0000000000
--- a/numba_dpex/tests/test_dpnp_functions.py
+++ /dev/null
@@ -1,45 +0,0 @@
-#! /usr/bin/env python
-
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import numpy as np
-import pytest
-from numba import njit
-
-from numba_dpex.tests._helper import (
-    assert_auto_offloading,
-    dpnp_debug,
-    skip_no_dpnp,
-    skip_no_opencl_gpu,
-)
-
-
-@skip_no_opencl_gpu
-@skip_no_dpnp
-class Testdpnp_functions:
-    N = 10
-
-    a = np.array(np.random.random(N), dtype=np.float32)
-    b = np.array(np.random.random(N), dtype=np.float32)
-    tys = [np.int32, np.uint32, np.int64, np.uint64, np.float32, np.double]
-
-    def test_dpnp_interacting_with_parfor(self):
-        def f(a, b):
-            c = np.sum(a)
-            e = np.add(b, a)
-            d = c + e
-            return d
-
-        device = dpctl.SyclDevice("opencl:gpu")
-        with dpctl.device_context(
-            device
-        ), assert_auto_offloading(), dpnp_debug():
-            njit_f = njit(f)
-            got = njit_f(self.a, self.b)
-        expected = f(self.a, self.b)
-
-        max_abs_err = got.sum() - expected.sum()
-        assert max_abs_err < 1e-4
diff --git a/numba_dpex/tests/test_dpnp_iface.py b/numba_dpex/tests/test_dpnp_iface.py
deleted file mode 100644
index d117dfe632..0000000000
--- a/numba_dpex/tests/test_dpnp_iface.py
+++ /dev/null
@@ -1,19 +0,0 @@
-#! /usr/bin/env python
-
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import pytest
-
-dpnp = pytest.importorskip("dpnp")
-
-
-def test_import_dpnp():
-    """Test that import dpnp works"""
-    import dpnp
-
-
-def test_import_dpnp_fptr_interface():
-    """Test that we can import dpnp_fptr_interface if dpnp is installed"""
-    from numba_dpex.dpnp_iface import dpnp_fptr_interface
diff --git a/numba_dpex/tests/test_offload_diagnostics.py b/numba_dpex/tests/test_offload_diagnostics.py
deleted file mode 100644
index a69cdc26a8..0000000000
--- a/numba_dpex/tests/test_offload_diagnostics.py
+++ /dev/null
@@ -1,38 +0,0 @@
-# Copyright 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import numpy as np
-from numba import njit, prange
-from numba.tests.support import captured_stdout
-
-import numba_dpex as dpex
-from numba_dpex import config as dpex_config
-from numba_dpex.tests._helper import skip_no_opencl_gpu
-
-
-@skip_no_opencl_gpu
-class TestOffloadDiagnostics:
-    def test_parfor(self):
-        def prange_func():
-            n = 10
-            a = np.ones((n), dtype=np.float64)
-            b = np.ones((n), dtype=np.float64)
-            c = np.ones((n), dtype=np.float64)
-            for i in prange(n // 2):
-                a[i] = b[i] + c[i]
-
-            return a
-
-        device = dpctl.SyclDevice("opencl:gpu")
-        with dpctl.device_context(device):
-            dpex_config.OFFLOAD_DIAGNOSTICS = 1
-            jitted = njit(parallel=True)(prange_func)
-
-            with captured_stdout() as got:
-                jitted()
-
-            dpex_config.OFFLOAD_DIAGNOSTICS = 0
-            assert "Auto-offloading" in got.getvalue()
-            assert "Device -" in got.getvalue()
diff --git a/numba_dpex/tests/test_parfor_lower_message.py b/numba_dpex/tests/test_parfor_lower_message.py
deleted file mode 100644
index 81ec974c12..0000000000
--- a/numba_dpex/tests/test_parfor_lower_message.py
+++ /dev/null
@@ -1,37 +0,0 @@
-# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import numpy as np
-from numba import njit, prange
-from numba.tests.support import captured_stdout
-
-from numba_dpex import config
-from numba_dpex.tests._helper import skip_no_opencl_gpu
-
-
-def prange_example():
-    n = 10
-    a = np.ones((n), dtype=np.float64)
-    b = np.ones((n), dtype=np.float64)
-    c = np.ones((n), dtype=np.float64)
-    for i in prange(n // 2):
-        a[i] = b[i] + c[i]
-
-    return a
-
-
-@skip_no_opencl_gpu
-class TestParforMessage:
-    def test_parfor_message(self):
-        device = dpctl.SyclDevice("opencl:gpu")
-        with dpctl.device_context(device):
-            config.DEBUG = 1
-            jitted = njit(prange_example)
-
-            with captured_stdout() as got:
-                jitted()
-
-            config.DEBUG = 0
-            assert "Parfor offloaded " in got.getvalue()
diff --git a/numba_dpex/tests/test_vectorize.py b/numba_dpex/tests/test_vectorize.py
index 00071c7e34..49cb358c53 100644
--- a/numba_dpex/tests/test_vectorize.py
+++ b/numba_dpex/tests/test_vectorize.py
@@ -9,7 +9,7 @@
 import pytest
 from numba import float32, float64, int32, int64, njit, vectorize
 
-from numba_dpex.tests._helper import assert_auto_offloading, filter_strings
+from numba_dpex.tests._helper import filter_strings
 
 list_of_shape = [
     (100, 100),
@@ -23,29 +23,6 @@ def shape(request):
     return request.param
 
 
-@pytest.mark.xfail
-@pytest.mark.parametrize("filter_str", filter_strings)
-def test_njit(filter_str):
-    @vectorize(nopython=True)
-    def axy(a, x, y):
-        return a * x + y
-
-    def f(a0, a1):
-        return np.cos(axy(a0, np.sin(a1) - 1.0, 1.0))
-
-    A = np.random.random(10)
-    B = np.random.random(10)
-
-    device = dpctl.SyclDevice(filter_str)
-    with dpctl.device_context(device), assert_auto_offloading():
-        f_njit = njit(f)
-        expected = f_njit(A, B)
-        actual = f(A, B)
-
-        max_abs_err = expected.sum() - actual.sum()
-        assert max_abs_err < 1e-5
-
-
 list_of_dtype = [
     (np.int32, int32),
     (np.float32, float32),
diff --git a/numba_dpex/tests/test_with_context.py b/numba_dpex/tests/test_with_context.py
deleted file mode 100644
index f86e166d93..0000000000
--- a/numba_dpex/tests/test_with_context.py
+++ /dev/null
@@ -1,99 +0,0 @@
-# Copyright 2020 - 2023 Intel Corporation
-#
-# SPDX-License-Identifier: Apache-2.0
-
-import dpctl
-import numpy as np
-import pytest
-from numba import njit
-from numba.tests.support import captured_stdout
-
-import numba_dpex as dpex
-from numba_dpex import config
-
-from ._helper import (
-    assert_auto_offloading,
-    filter_strings,
-    skip_no_opencl_cpu,
-    skip_no_opencl_gpu,
-)
-
-
-def scenario(filter_str, context):
-    @njit
-    def func(a, b):
-        return a + b
-
-    a = np.ones((64), dtype=np.int32)
-    b = np.ones((64), dtype=np.int32)
-
-    device = dpctl.SyclDevice(filter_str)
-    with context(device):
-        func(a, b)
-
-
-@pytest.mark.parametrize("filter_str", filter_strings)
-@pytest.mark.parametrize(
-    "context",
-    [
-        dpex.offload_to_sycl_device,
-        dpctl.device_context,
-    ],
-)
-def test_dpctl_device_context_affects_numba_pipeline(filter_str, context):
-    with assert_auto_offloading():
-        scenario(filter_str, context)
-
-
-class TestWithDeviceContext:
-    @skip_no_opencl_gpu
-    def test_with_device_context_gpu(self):
-        @njit
-        def nested_func(a, b):
-            np.sin(a, b)
-
-        @njit
-        def func(b):
-            a = np.ones((64), dtype=np.float64)
-            nested_func(a, b)
-
-        config.DEBUG = 1
-        expected = np.ones((64), dtype=np.float64)
-        got_gpu = np.ones((64), dtype=np.float64)
-
-        with captured_stdout() as got_gpu_message:
-            device = dpctl.SyclDevice("opencl:gpu")
-            with dpctl.device_context(device):
-                func(got_gpu)
-
-        config.DEBUG = 0
-        func(expected)
-
-        np.testing.assert_array_equal(expected, got_gpu)
-        assert "Parfor offloaded to opencl:gpu" in got_gpu_message.getvalue()
-
-    @skip_no_opencl_cpu
-    def test_with_device_context_cpu(self):
-        @njit
-        def nested_func(a, b):
-            np.sin(a, b)
-
-        @njit
-        def func(b):
-            a = np.ones((64), dtype=np.float64)
-            nested_func(a, b)
-
-        config.DEBUG = 1
-        expected = np.ones((64), dtype=np.float64)
-        got_cpu = np.ones((64), dtype=np.float64)
-
-        with captured_stdout() as got_cpu_message:
-            device = dpctl.SyclDevice("opencl:cpu")
-            with dpctl.device_context(device):
-                func(got_cpu)
-
-        config.DEBUG = 0
-        func(expected)
-
-        np.testing.assert_array_equal(expected, got_cpu)
-        assert "Parfor offloaded to opencl:cpu" in got_cpu_message.getvalue()