From b5b70b35cdbf6767fd2d1b8af147756c537d6cef Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Thu, 4 May 2023 02:12:43 -0500 Subject: [PATCH] Big spring cleaning --- numba_dpex/__init__.py | 4 +- numba_dpex/core/compiler.py | 1 - numba_dpex/core/offload_dispatcher.py | 60 - .../passes/rename_numpy_functions_pass.py | 361 ----- numba_dpex/core/pipelines/offload_compiler.py | 204 --- numba_dpex/dpnp_iface/dpnp_array_ops_impl.py | 344 ----- numba_dpex/dpnp_iface/dpnp_fptr_interface.pyx | 267 ---- numba_dpex/dpnp_iface/dpnp_indexing.py | 115 -- numba_dpex/dpnp_iface/dpnp_linalgimpl.py | 734 --------- numba_dpex/dpnp_iface/dpnp_logic.py | 72 - numba_dpex/dpnp_iface/dpnp_manipulation.py | 82 - numba_dpex/dpnp_iface/dpnp_randomimpl.py | 1374 ----------------- .../dpnp_iface/dpnp_sort_search_countimpl.py | 266 ---- numba_dpex/dpnp_iface/dpnp_statisticsimpl.py | 388 ----- numba_dpex/dpnp_iface/dpnp_stubs_impl.py | 79 - .../dpnp_iface/dpnp_transcendentalsimpl.py | 194 --- numba_dpex/dpnp_iface/dpnpdecl.py | 45 - numba_dpex/retarget.py | 61 - numba_dpex/tests/_helper.py | 17 - .../passes/test_rename_numpy_function_pass.py | 113 -- .../kernel_tests/test_compute_follows_data.py | 238 --- numba_dpex/tests/njit_tests/__init__.py | 6 - numba_dpex/tests/njit_tests/dpnp/__init__.py | 5 - numba_dpex/tests/njit_tests/dpnp/_helper.py | 23 - .../njit_tests/dpnp/test_numpy_array_ops.py | 142 -- .../njit_tests/dpnp/test_numpy_indexing.py | 63 - .../njit_tests/dpnp/test_numpy_linalg.py | 326 ---- .../tests/njit_tests/dpnp/test_numpy_logic.py | 48 - .../dpnp/test_numpy_manipulation.py | 33 - .../tests/njit_tests/dpnp/test_numpy_rng.py | 229 --- .../dpnp/test_numpy_sort_search_count.py | 117 -- .../njit_tests/dpnp/test_numpy_statistics.py | 86 -- .../dpnp/test_numpy_transcendentals.py | 138 -- .../tests/test_controllable_fallback.py | 75 - numba_dpex/tests/test_debuginfo.py | 200 --- numba_dpex/tests/test_dpnp_functions.py | 45 - numba_dpex/tests/test_dpnp_iface.py | 19 - numba_dpex/tests/test_offload_diagnostics.py | 38 - numba_dpex/tests/test_parfor_lower_message.py | 37 - numba_dpex/tests/test_with_context.py | 99 -- setup.py | 14 - 41 files changed, 1 insertion(+), 6761 deletions(-) delete mode 100644 numba_dpex/core/offload_dispatcher.py delete mode 100644 numba_dpex/core/passes/rename_numpy_functions_pass.py delete mode 100644 numba_dpex/core/pipelines/offload_compiler.py delete mode 100644 numba_dpex/dpnp_iface/dpnp_array_ops_impl.py delete mode 100644 numba_dpex/dpnp_iface/dpnp_fptr_interface.pyx delete mode 100644 numba_dpex/dpnp_iface/dpnp_indexing.py delete mode 100644 numba_dpex/dpnp_iface/dpnp_linalgimpl.py delete mode 100644 numba_dpex/dpnp_iface/dpnp_logic.py delete mode 100644 numba_dpex/dpnp_iface/dpnp_manipulation.py delete mode 100644 numba_dpex/dpnp_iface/dpnp_randomimpl.py delete mode 100644 numba_dpex/dpnp_iface/dpnp_sort_search_countimpl.py delete mode 100644 numba_dpex/dpnp_iface/dpnp_statisticsimpl.py delete mode 100644 numba_dpex/dpnp_iface/dpnp_stubs_impl.py delete mode 100644 numba_dpex/dpnp_iface/dpnp_transcendentalsimpl.py delete mode 100644 numba_dpex/dpnp_iface/dpnpdecl.py delete mode 100644 numba_dpex/retarget.py delete mode 100644 numba_dpex/tests/core/passes/test_rename_numpy_function_pass.py delete mode 100644 numba_dpex/tests/kernel_tests/test_compute_follows_data.py delete mode 100644 numba_dpex/tests/njit_tests/__init__.py delete mode 100644 numba_dpex/tests/njit_tests/dpnp/__init__.py delete mode 100644 numba_dpex/tests/njit_tests/dpnp/_helper.py delete mode 100644 numba_dpex/tests/njit_tests/dpnp/test_numpy_array_ops.py delete mode 100644 numba_dpex/tests/njit_tests/dpnp/test_numpy_indexing.py delete mode 100644 numba_dpex/tests/njit_tests/dpnp/test_numpy_linalg.py delete mode 100644 numba_dpex/tests/njit_tests/dpnp/test_numpy_logic.py delete mode 100644 numba_dpex/tests/njit_tests/dpnp/test_numpy_manipulation.py delete mode 100644 numba_dpex/tests/njit_tests/dpnp/test_numpy_rng.py delete mode 100644 numba_dpex/tests/njit_tests/dpnp/test_numpy_sort_search_count.py delete mode 100644 numba_dpex/tests/njit_tests/dpnp/test_numpy_statistics.py delete mode 100644 numba_dpex/tests/njit_tests/dpnp/test_numpy_transcendentals.py delete mode 100644 numba_dpex/tests/test_controllable_fallback.py delete mode 100644 numba_dpex/tests/test_debuginfo.py delete mode 100644 numba_dpex/tests/test_dpnp_functions.py delete mode 100644 numba_dpex/tests/test_dpnp_iface.py delete mode 100644 numba_dpex/tests/test_offload_diagnostics.py delete mode 100644 numba_dpex/tests/test_parfor_lower_message.py delete mode 100644 numba_dpex/tests/test_with_context.py diff --git a/numba_dpex/__init__.py b/numba_dpex/__init__.py index d043aea5ce..fc1adecfa4 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,7 +110,6 @@ 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 @@ -150,4 +148,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..2be71324a5 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 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: ) ['$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: ) ['$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/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 = kernel_data.ptr - - return 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/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/numba_dpex/tests/_helper.py b/numba_dpex/tests/_helper.py index 515a4a31d8..978a7b1327 100644 --- a/numba_dpex/tests/_helper.py +++ b/numba_dpex/tests/_helper.py @@ -141,23 +141,6 @@ 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 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/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() 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_debuginfo.py b/numba_dpex/tests/test_debuginfo.py deleted file mode 100644 index 679bf53579..0000000000 --- a/numba_dpex/tests/test_debuginfo.py +++ /dev/null @@ -1,200 +0,0 @@ -#! /usr/bin/env python - -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import re - -import pytest - -import numba_dpex as dpex -from numba_dpex import float32, int32, usm_ndarray -from numba_dpex.core.descriptor import dpex_kernel_target -from numba_dpex.tests._helper import override_config - -debug_options = [True, False] - -f32arrty = usm_ndarray(ndim=1, dtype=float32, layout="C") - - -@pytest.fixture(params=debug_options) -def debug_option(request): - return request.param - - -def get_kernel_ir(fn, sig, debug=False): - kernel = dpex.core.kernel_interface.spirv_kernel.SpirvKernel( - fn, fn.__name__ - ) - kernel.compile( - args=sig, - target_ctx=dpex_kernel_target.target_context, - typing_ctx=dpex_kernel_target.typing_context, - debug=debug, - compile_flags=None, - ) - return kernel.llvm_module - - -def make_check(ir, val_to_search): - """ - Check the compiled assembly for debuginfo. - """ - - m = re.search(val_to_search, ir, re.I) - got = m is not None - return got - - -def test_debug_flag_generates_ir_with_debuginfo(debug_option): - """ - Check debug info is emitting to IR if debug parameter is set to True - """ - - def foo(x): - x = 1 # noqa - - sig = (int32,) - kernel_ir = get_kernel_ir(foo, sig, debug=debug_option) - tag = "!dbg" - - if debug_option: - assert tag in kernel_ir - else: - assert tag not in kernel_ir - - -def test_debug_info_locals_vars_on_no_opt(): - """ - Check llvm debug tag DILocalVariable is emitting to IR for all variables - if debug parameter is set to True and optimization is O0 - """ - - def foo(var_a, var_b, var_c): - i = dpex.get_global_id(0) - var_c[i] = var_a[i] + var_b[i] - - ir_tags = [ - '!DILocalVariable(name: "var_a"', - '!DILocalVariable(name: "var_b"', - '!DILocalVariable(name: "var_c"', - '!DILocalVariable(name: "i"', - ] - sig = (f32arrty, f32arrty, f32arrty) - - with override_config("OPT", 0): - kernel_ir = get_kernel_ir(foo, sig, debug=True) - - for tag in ir_tags: - assert tag in kernel_ir - - -def test_debug_kernel_local_vars_in_ir(): - """ - Check llvm debug tag DILocalVariable is emitting to IR for variables - created in kernel - """ - - def foo(arr): - index = dpex.get_global_id(0) - local_d = 9 * 99 + 5 - arr[index] = local_d + 100 - - ir_tags = [ - '!DILocalVariable(name: "index"', - '!DILocalVariable(name: "local_d"', - ] - sig = (f32arrty,) - kernel_ir = get_kernel_ir(foo, sig, debug=True) - - for tag in ir_tags: - assert tag in kernel_ir - - -def test_debug_flag_generates_ir_with_debuginfo_for_func(debug_option): - """ - Check debug info is emitting to IR if debug parameter is set to True - """ - - @dpex.func(debug=debug_option) - def func_sum(a, b): - result = a + b - return result - - def data_parallel_sum(a, b, c): - i = dpex.get_global_id(0) - c[i] = func_sum(a[i], b[i]) - - ir_tags = [ - r'\!DISubprogram\(name: ".*func_sum\$?\d*"', - r'\!DISubprogram\(name: ".*data_parallel_sum\$?\d*"', - ] - - sig = (f32arrty, f32arrty, f32arrty) - - kernel_ir = get_kernel_ir(data_parallel_sum, sig, debug=debug_option) - - for tag in ir_tags: - assert debug_option == make_check(kernel_ir, tag) - - -def test_env_var_generates_ir_with_debuginfo_for_func(debug_option): - """ - Check debug info is emitting to IR if NUMBA_DPEX_DEBUGINFO is set to 1 - """ - - @dpex.func - def func_sum(a, b): - result = a + b - return result - - def data_parallel_sum(a, b, c): - i = dpex.get_global_id(0) - c[i] = func_sum(a[i], b[i]) - - ir_tags = [ - r'\!DISubprogram\(name: ".*func_sum\$?\d*"', - r'\!DISubprogram\(name: ".*data_parallel_sum\$\d*"', - ] - - sig = (f32arrty, f32arrty, f32arrty) - - with override_config("DEBUGINFO_DEFAULT", int(debug_option)): - kernel_ir = get_kernel_ir(data_parallel_sum, sig) - - for tag in ir_tags: - assert debug_option == make_check(kernel_ir, tag) - - -def test_debuginfo_DISubprogram_linkageName(): - def func(a, b): - i = dpex.get_global_id(0) - b[i] = a[i] - - ir_tags = [ - r'\!DISubprogram\(.*linkageName: ".*func.*"', - ] - - sig = (f32arrty, f32arrty) - kernel_ir = get_kernel_ir(func, sig, debug=True) - - for tag in ir_tags: - assert make_check(kernel_ir, tag) - - -def test_debuginfo_DICompileUnit_language_and_producer(): - def func(a, b): - i = dpex.get_global_id(0) - b[i] = a[i] - - ir_tags = [ - r"\!DICompileUnit\(language: DW_LANG_C_plus_plus,", - ] - - sig = (f32arrty, f32arrty) - - kernel_ir = get_kernel_ir(func, sig, debug=True) - - for tag in ir_tags: - assert make_check(kernel_ir, tag) 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_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() 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=[