diff --git a/numba_cuda/numba/cuda/__init__.py b/numba_cuda/numba/cuda/__init__.py index ad92b20ac..d0ff4ba55 100644 --- a/numba_cuda/numba/cuda/__init__.py +++ b/numba_cuda/numba/cuda/__init__.py @@ -7,12 +7,14 @@ import warnings import sys + # Re-export types itself import numba.cuda.types as types # Re-export all type names from numba.cuda.types import * +HAS_NUMBA = importlib.util.find_spec("numba") is not None # Require NVIDIA CUDA bindings at import time if not ( diff --git a/numba_cuda/numba/cuda/compiler.py b/numba_cuda/numba/cuda/compiler.py index cd9b6b76e..a1653c192 100644 --- a/numba_cuda/numba/cuda/compiler.py +++ b/numba_cuda/numba/cuda/compiler.py @@ -446,7 +446,7 @@ def run_pass(self, state): @register_pass(mutates_CFG=True, analysis_only=False) class CUDANativeLowering(BaseNativeLowering): """Lowering pass for a CUDA native function IR described solely in terms of - Numba's standard `numba.core.ir` nodes.""" + Numba's standard `numba.cuda.core.ir` nodes.""" _name = "cuda_native_lowering" diff --git a/numba_cuda/numba/cuda/core/analysis.py b/numba_cuda/numba/cuda/core/analysis.py index 4cc19cc37..bda7fdf2f 100644 --- a/numba_cuda/numba/cuda/core/analysis.py +++ b/numba_cuda/numba/cuda/core/analysis.py @@ -4,13 +4,13 @@ from collections import namedtuple, defaultdict from numba.cuda import types from numba.cuda.core import ir -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.core import consts import operator from functools import reduce from .controlflow import CFGraph -from numba.misc import special +from numba.cuda.misc import special # # Analysis related to variable lifetime @@ -354,7 +354,7 @@ def rewrite_tuple_len(val, func_ir, called_args): if isinstance(argty, types.BaseTuple): rewrite_statement(func_ir, stmt, argty.count) - from numba.core.ir_utils import get_definition, guard + from numba.cuda.core.ir_utils import get_definition, guard for blk in func_ir.blocks.values(): for stmt in blk.body: diff --git a/numba_cuda/numba/cuda/core/base.py b/numba_cuda/numba/cuda/core/base.py index 250df455e..c20b50e30 100644 --- a/numba_cuda/numba/cuda/core/base.py +++ b/numba_cuda/numba/cuda/core/base.py @@ -14,7 +14,7 @@ from numba.cuda.core import imputils, targetconfig, funcdesc from numba.cuda import cgutils, debuginfo, types, utils, datamodel, config -from numba.core import errors +from numba.cuda.core import errors from numba.core.compiler_lock import global_compiler_lock from numba.cuda.core.pythonapi import PythonAPI from numba.cuda.core.imputils import ( diff --git a/numba_cuda/numba/cuda/core/bytecode.py b/numba_cuda/numba/cuda/core/bytecode.py index cda186133..d86b6b496 100644 --- a/numba_cuda/numba/cuda/core/bytecode.py +++ b/numba_cuda/numba/cuda/core/bytecode.py @@ -8,8 +8,8 @@ from types import CodeType, ModuleType -from numba.core import errors -from numba.core import serialize +from numba.cuda.core import errors +from numba.cuda import serialize from numba.cuda import utils from numba.cuda.utils import PYVERSION diff --git a/numba_cuda/numba/cuda/core/byteflow.py b/numba_cuda/numba/cuda/core/byteflow.py index 4be6d4c10..d31ed7dd8 100644 --- a/numba_cuda/numba/cuda/core/byteflow.py +++ b/numba_cuda/numba/cuda/core/byteflow.py @@ -17,7 +17,7 @@ _lazy_pformat, ) from numba.cuda.core.controlflow import NEW_BLOCKERS, CFGraph -from numba.core.ir import Loc +from numba.cuda.core.ir import Loc from numba.cuda.errors import UnsupportedBytecodeError diff --git a/numba_cuda/numba/cuda/core/compiler.py b/numba_cuda/numba/cuda/core/compiler.py index cba5b43dc..9cb8a1840 100644 --- a/numba_cuda/numba/cuda/core/compiler.py +++ b/numba_cuda/numba/cuda/core/compiler.py @@ -3,7 +3,7 @@ from numba.cuda.core.tracing import event -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.core.errors import CompilerError from numba.cuda.core import callconv, config, bytecode diff --git a/numba_cuda/numba/cuda/core/compiler_machinery.py b/numba_cuda/numba/cuda/core/compiler_machinery.py index bb334f71a..cb91f057e 100644 --- a/numba_cuda/numba/cuda/core/compiler_machinery.py +++ b/numba_cuda/numba/cuda/core/compiler_machinery.py @@ -8,7 +8,7 @@ from numba.core.compiler_lock import global_compiler_lock -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.core import config from numba.cuda import utils from numba.cuda.core import transforms diff --git a/numba_cuda/numba/cuda/core/config.py b/numba_cuda/numba/cuda/core/config.py index 92db61688..5774c8b79 100644 --- a/numba_cuda/numba/cuda/core/config.py +++ b/numba_cuda/numba/cuda/core/config.py @@ -212,11 +212,6 @@ def _readenv(name, ctor, default): def optional_str(x): return str(x) if x is not None else None - # Type casting rules selection - USE_LEGACY_TYPE_SYSTEM = _readenv( - "NUMBA_USE_LEGACY_TYPE_SYSTEM", int, 1 - ) - # developer mode produces full tracebacks, disables help instructions DEVELOPER_MODE = _readenv("NUMBA_DEVELOPER_MODE", int, 0) diff --git a/numba_cuda/numba/cuda/core/controlflow.py b/numba_cuda/numba/cuda/core/controlflow.py index 528aa3678..4741cc1b2 100644 --- a/numba_cuda/numba/cuda/core/controlflow.py +++ b/numba_cuda/numba/cuda/core/controlflow.py @@ -5,7 +5,7 @@ import functools import sys -from numba.core.ir import Loc +from numba.cuda.core.ir import Loc from numba.cuda.core.errors import UnsupportedError from numba.cuda.utils import PYVERSION diff --git a/numba_cuda/numba/cuda/core/entrypoints.py b/numba_cuda/numba/cuda/core/entrypoints.py index fb8a6e06f..13679556f 100644 --- a/numba_cuda/numba/cuda/core/entrypoints.py +++ b/numba_cuda/numba/cuda/core/entrypoints.py @@ -5,7 +5,7 @@ import warnings from importlib import metadata as importlib_metadata - +from numba.cuda import HAS_NUMBA _already_initialized = False logger = logging.getLogger(__name__) @@ -16,12 +16,10 @@ def init_all(): If extensions have already been initialized, this function does nothing. """ - try: + if HAS_NUMBA: from numba.core import entrypoints entrypoints.init_all() - except ImportError: - pass global _already_initialized if _already_initialized: diff --git a/numba_cuda/numba/cuda/core/errors.py b/numba_cuda/numba/cuda/core/errors.py index aef7b26e5..0d15e3c6e 100644 --- a/numba_cuda/numba/cuda/core/errors.py +++ b/numba_cuda/numba/cuda/core/errors.py @@ -9,12 +9,14 @@ import contextlib import os import warnings -import numba.core.config +import numba.cuda.core.config import numpy as np from collections import defaultdict from functools import wraps from abc import abstractmethod +import numba_cuda + # Filled at the end __all__ = [] @@ -24,7 +26,7 @@ def _is_numba_core_config_loaded(): To detect if numba.core.config has been initialized due to circular imports. """ try: - numba.core.config + numba.cuda.core.config except AttributeError: return False else: @@ -454,7 +456,7 @@ def termcolor(): This should not have happened, a problem has occurred in Numba's internals. You are currently using Numba version %s. %s -""" % (numba.__version__, feedback_details) +""" % (numba_cuda.__version__, feedback_details) error_extras = dict() error_extras["unsupported_error"] = unsupported_error_info @@ -813,7 +815,7 @@ def __init__(self, arg_indices, fold_arguments=None, loc=None): def bind_fold_arguments(self, fold_arguments): """Bind the fold_arguments function""" # to avoid circular import - from numba.core.utils import chain_exception + from numba.cuda.core.utils import chain_exception e = ForceLiteralArg(self.requested_args, fold_arguments, loc=self.loc) return chain_exception(e, self) diff --git a/numba_cuda/numba/cuda/core/event.py b/numba_cuda/numba/cuda/core/event.py index f5bb27a5d..29b92da6c 100644 --- a/numba_cuda/numba/cuda/core/event.py +++ b/numba_cuda/numba/cuda/core/event.py @@ -43,7 +43,7 @@ from timeit import default_timer as timer from contextlib import contextmanager, ExitStack from collections import defaultdict -from numba.core import config +from numba.cuda.core import config from numba.cuda import utils diff --git a/numba_cuda/numba/cuda/core/imputils.py b/numba_cuda/numba/cuda/core/imputils.py index fdd3f22fd..2881c2771 100644 --- a/numba_cuda/numba/cuda/core/imputils.py +++ b/numba_cuda/numba/cuda/core/imputils.py @@ -492,7 +492,7 @@ def force_error_model(context, model_name="numpy"): """ Temporarily change the context's error model. """ - from numba.core import callconv + from numba.cuda.core import callconv old_error_model = context.error_model context.error_model = callconv.create_error_model(model_name, context) diff --git a/numba_cuda/numba/cuda/core/inline_closurecall.py b/numba_cuda/numba/cuda/core/inline_closurecall.py index c45f3265a..4b5dd09f7 100644 --- a/numba_cuda/numba/cuda/core/inline_closurecall.py +++ b/numba_cuda/numba/cuda/core/inline_closurecall.py @@ -5,9 +5,10 @@ import copy import ctypes import numba.cuda.core.analysis +from numba.cuda import HAS_NUMBA from numba.cuda import types, config, cgutils from numba.cuda.core import ir -from numba.core import errors +from numba.cuda.core import errors from numba.cuda import typing, utils from numba.cuda.core.ir_utils import ( next_label, @@ -43,7 +44,6 @@ from numba.cuda.np.unsafe.ndarray import empty_inferred as unsafe_empty_inferred import numpy as np import operator -from numba.cuda.misc.special import prange """ Variable enable_inline_arraycall is only used for testing purpose. @@ -85,7 +85,6 @@ def __init__(self, func_ir, parallel_options, swapped=None, typed=False): self.func_ir = func_ir self.parallel_options = parallel_options self.swapped = swapped - self._processed_stencils = [] self.typed = typed def run(self): @@ -126,11 +125,6 @@ def run(self): modified = True break # because block structure changed - if guard( - self._inline_stencil, instr, call_name, func_def - ): - modified = True - if enable_inline_arraycall: # Identify loop structure if modified: @@ -216,91 +210,6 @@ def reduce_func(f, A, v=None): ) return True - def _inline_stencil(self, instr, call_name, func_def): - from numba.stencils.stencil import StencilFunc - - lhs = instr.target - expr = instr.value - # We keep the escaping variables of the stencil kernel - # alive by adding them to the actual kernel call as extra - # keyword arguments, which is ignored anyway. - if ( - isinstance(func_def, ir.Global) - and func_def.name == "stencil" - and isinstance(func_def.value, StencilFunc) - ): - if expr.kws: - expr.kws += func_def.value.kws - else: - expr.kws = func_def.value.kws - return True - # Otherwise we proceed to check if it is a call to numba.stencil - require( - call_name == ("stencil", "numba.stencils.stencil") - or call_name == ("stencil", "numba") - ) - require(expr not in self._processed_stencils) - self._processed_stencils.append(expr) - if not len(expr.args) == 1: - raise ValueError( - "As a minimum Stencil requires a kernel as an argument" - ) - stencil_def = guard(get_definition, self.func_ir, expr.args[0]) - require( - isinstance(stencil_def, ir.Expr) - and stencil_def.op == "make_function" - ) - kernel_ir = get_ir_of_code( - self.func_ir.func_id.func.__globals__, stencil_def.code - ) - options = dict(expr.kws) - if "neighborhood" in options: - fixed = guard(self._fix_stencil_neighborhood, options) - if not fixed: - raise ValueError( - "stencil neighborhood option should be a tuple" - " with constant structure such as ((-w, w),)" - ) - if "index_offsets" in options: - fixed = guard(self._fix_stencil_index_offsets, options) - if not fixed: - raise ValueError( - "stencil index_offsets option should be a tuple" - " with constant structure such as (offset, )" - ) - sf = StencilFunc(kernel_ir, "constant", options) - sf.kws = expr.kws # hack to keep variables live - sf_global = ir.Global("stencil", sf, expr.loc) - self.func_ir._definitions[lhs.name] = [sf_global] - instr.value = sf_global - return True - - def _fix_stencil_neighborhood(self, options): - """ - Extract the two-level tuple representing the stencil neighborhood - from the program IR to provide a tuple to StencilFunc. - """ - # build_tuple node with neighborhood for each dimension - dims_build_tuple = get_definition(self.func_ir, options["neighborhood"]) - require(hasattr(dims_build_tuple, "items")) - res = [] - for window_var in dims_build_tuple.items: - win_build_tuple = get_definition(self.func_ir, window_var) - require(hasattr(win_build_tuple, "items")) - res.append(tuple(win_build_tuple.items)) - options["neighborhood"] = tuple(res) - return True - - def _fix_stencil_index_offsets(self, options): - """ - Extract the tuple representing the stencil index offsets - from the program IR to provide to StencilFunc. - """ - offset_tuple = get_definition(self.func_ir, options["index_offsets"]) - require(hasattr(offset_tuple, "items")) - options["index_offsets"] = tuple(offset_tuple.items) - return True - def _inline_closure(self, work_list, block, i, func_def): require( isinstance(func_def, ir.Expr) and func_def.op == "make_function" @@ -327,13 +236,11 @@ def check_reduce_func(func_ir, func_var): analysis" ) if isinstance(reduce_func, (ir.FreeVar, ir.Global)): - try: + if HAS_NUMBA: from numba.core.registry import CPUDispatcher if not isinstance(reduce_func.value, CPUDispatcher): raise ValueError("Invalid reduction function") - except ImportError: - pass # pull out the python function for inlining reduce_func = reduce_func.value.py_func @@ -1060,10 +967,7 @@ def _find_iter_range(func_ir, range_iter_var, swapped): func_var = range_def.func func_def = get_definition(func_ir, func_var) debug_print("func_var = ", func_var, " func_def = ", func_def) - require( - isinstance(func_def, ir.Global) - and (func_def.value is range or func_def.value == prange) - ) + require(isinstance(func_def, ir.Global) and func_def.value is range) nargs = len(range_def.args) swapping = [('"array comprehension"', "closure of"), range_def.func.loc] if nargs == 1: @@ -1138,17 +1042,6 @@ def codegen(context, builder, sig, args): count_const = intp_t(tuplety.count) return impl_ret_untracked(context, builder, intp_t, count_const) - return signature(types.intp, val), codegen - elif isinstance(val, types.ListTypeIteratorType): - - def codegen(context, builder, sig, args): - (value,) = args - intp_t = context.get_value_type(types.intp) - from numba.typed.listobject import ListIterInstance - - iterobj = ListIterInstance(context, builder, sig.args[0], value) - return impl_ret_untracked(context, builder, intp_t, iterobj.size) - return signature(types.intp, val), codegen else: msg = ( diff --git a/numba_cuda/numba/cuda/core/interpreter.py b/numba_cuda/numba/cuda/core/interpreter.py index 4e8fefac0..d63f6cf11 100644 --- a/numba_cuda/numba/cuda/core/interpreter.py +++ b/numba_cuda/numba/cuda/core/interpreter.py @@ -8,7 +8,7 @@ import logging import textwrap -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.core import config, ir from numba.cuda.errors import UnsupportedBytecodeError from numba.cuda.core.errors import ( diff --git a/numba_cuda/numba/cuda/core/ir.py b/numba_cuda/numba/cuda/core/ir.py index 9951fdfbe..807733fea 100644 --- a/numba_cuda/numba/cuda/core/ir.py +++ b/numba_cuda/numba/cuda/core/ir.py @@ -1,7 +1,6 @@ # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: BSD-2-Clause -import numba from collections import defaultdict import copy import itertools @@ -15,7 +14,11 @@ from functools import total_ordering from io import StringIO -from numba.core import errors +from numba.cuda import HAS_NUMBA + +if HAS_NUMBA: + import numba +from numba.cuda.core import errors from numba.cuda.core import config from numba.cuda.utils import UNARY_BUILTINS_TO_OPERATORS, OPERATORS_TO_BUILTINS from numba.cuda.core.errors import ( @@ -776,7 +779,10 @@ def __repr__(self): class Del(Stmt): def __init__(self, value, loc): assert isinstance(value, str) - assert isinstance(loc, (Loc, numba.core.ir.Loc)) + if HAS_NUMBA: + assert isinstance(loc, (Loc, numba.core.ir.Loc)) + else: + assert isinstance(loc, (Loc)) self.value = value self.loc = loc @@ -1342,8 +1348,12 @@ class Block(EqualityCheckMixin): """A code block""" def __init__(self, scope, loc): - assert isinstance(scope, (Scope, numba.core.ir.Scope)) - assert isinstance(loc, (Loc, numba.core.ir.Loc)) + if HAS_NUMBA: + assert isinstance(scope, (Scope, numba.core.ir.Scope)) + assert isinstance(loc, (Loc, numba.core.ir.Loc)) + else: + assert isinstance(scope, Scope) + assert isinstance(loc, Loc) self.scope = scope self.body = [] self.loc = loc @@ -1708,8 +1718,8 @@ def dump(self, file=None): raise ValueError(msg) else: from pygments import highlight - from numba.misc.dump_style import NumbaIRLexer as lexer - from numba.misc.dump_style import by_colorscheme + from numba.cuda.misc.dump_style import NumbaIRLexer as lexer + from numba.cuda.misc.dump_style import by_colorscheme from pygments.formatters import Terminal256Formatter print( diff --git a/numba_cuda/numba/cuda/core/ir_utils.py b/numba_cuda/numba/cuda/core/ir_utils.py index c290cda12..225346985 100644 --- a/numba_cuda/numba/cuda/core/ir_utils.py +++ b/numba_cuda/numba/cuda/core/ir_utils.py @@ -9,7 +9,8 @@ import collections import warnings -import numba +import numba.cuda +from numba.cuda import HAS_NUMBA from numba.cuda import types from numba.cuda.core import ir from numba.cuda import typing @@ -816,10 +817,8 @@ def has_no_side_effect(rhs, lives, call_table): if ( call_list == ["empty", numpy] or call_list == [slice] - or call_list == ["stencil", numba] or call_list == ["log", numpy] or call_list == ["dtype", numpy] - or call_list == ["pndindex", numba] or call_list == ["ceil", math] or call_list == [max] or call_list == [int] @@ -831,7 +830,7 @@ def has_no_side_effect(rhs, lives, call_table): ): return True - try: + if HAS_NUMBA: from numba.core.registry import CPUDispatcher from numba.cuda.np.linalg import dot_3_mv_check_args @@ -839,8 +838,6 @@ def has_no_side_effect(rhs, lives, call_table): py_func = call_list[0].py_func if py_func == dot_3_mv_check_args: return True - except ImportError: - pass for f in remove_call_handlers: if f(rhs, lives, call_list): @@ -1992,7 +1989,7 @@ def __init__(self, f_ir): state = DummyPipeline(ir).state rewrites.rewrite_registry.apply("before-inference", state) - # call inline pass to handle cases like stencils and comprehensions + # call inline pass to handle cases like comprehensions swapped = {} # TODO: get this from diagnostics store from numba.cuda.core.inline_closurecall import InlineClosureCallPass @@ -2333,7 +2330,7 @@ def raise_on_unsupported_feature(func_ir, typemap): # check global function found = False if isinstance(val, pytypes.FunctionType): - found = val in {numba.gdb, numba.gdb_init} + found = val in {numba.cuda.gdb, numba.cuda.gdb_init} if not found: # freevar bind to intrinsic found = getattr(val, "_name", "") == "gdb_internal" if found: @@ -2491,7 +2488,7 @@ def legalize_single_scope(blocks): return len({blk.scope for blk in blocks.values()}) == 1 -def check_and_legalize_ir(func_ir, flags: "numba.core.flags.Flags"): +def check_and_legalize_ir(func_ir, flags: "numba.cuda.flags.Flags"): """ This checks that the IR presented is legal """ diff --git a/numba_cuda/numba/cuda/core/pythonapi.py b/numba_cuda/numba/cuda/core/pythonapi.py index f7b5e03b6..5a3fda59f 100644 --- a/numba_cuda/numba/cuda/core/pythonapi.py +++ b/numba_cuda/numba/cuda/core/pythonapi.py @@ -11,7 +11,7 @@ import ctypes from numba.cuda.cext import _helperlib -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.core import imputils from numba.cuda.utils import PYVERSION from numba.cuda import config, types, lowering, cgutils, serialize diff --git a/numba_cuda/numba/cuda/core/rewrites/static_raise.py b/numba_cuda/numba/cuda/core/rewrites/static_raise.py index 9860bc154..5dbac7834 100644 --- a/numba_cuda/numba/cuda/core/rewrites/static_raise.py +++ b/numba_cuda/numba/cuda/core/rewrites/static_raise.py @@ -1,7 +1,7 @@ # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: BSD-2-Clause -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.core import ir from numba.cuda.core import consts from numba.cuda.core.rewrites import register_rewrite, Rewrite diff --git a/numba_cuda/numba/cuda/core/sigutils.py b/numba_cuda/numba/cuda/core/sigutils.py index 04fa54c45..cd7f9a65e 100644 --- a/numba_cuda/numba/cuda/core/sigutils.py +++ b/numba_cuda/numba/cuda/core/sigutils.py @@ -1,15 +1,11 @@ # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: BSD-2-Clause -from numba.cuda import types, typing +from numba.cuda import types, typing, HAS_NUMBA -try: +if HAS_NUMBA: from numba.core.typing import Signature as CoreSignature - numba_sig_present = True -except ImportError: - numba_sig_present = False - def is_signature(sig): """ @@ -17,7 +13,7 @@ def is_signature(sig): specification (for user-facing APIs). """ sig_types = (str, tuple, typing.Signature) - if numba_sig_present: + if HAS_NUMBA: sig_types = (str, tuple, typing.Signature, CoreSignature) return isinstance(sig, sig_types) @@ -46,7 +42,7 @@ def normalize_signature(sig): args, return_type = parsed, None else: sig_types = (typing.Signature,) - if numba_sig_present: + if HAS_NUMBA: sig_types = (typing.Signature, CoreSignature) if isinstance(parsed, sig_types): args, return_type = parsed.args, parsed.return_type diff --git a/numba_cuda/numba/cuda/core/ssa.py b/numba_cuda/numba/cuda/core/ssa.py index 0deeae124..a157e83d8 100644 --- a/numba_cuda/numba/cuda/core/ssa.py +++ b/numba_cuda/numba/cuda/core/ssa.py @@ -19,7 +19,7 @@ from numba.cuda import config from numba.cuda.core import ir -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.core import ir_utils from numba.cuda.utils import _lazy_pformat from numba.cuda.core.analysis import compute_cfg_from_blocks diff --git a/numba_cuda/numba/cuda/core/targetconfig.py b/numba_cuda/numba/cuda/core/targetconfig.py index ac31fcd28..0559a130b 100644 --- a/numba_cuda/numba/cuda/core/targetconfig.py +++ b/numba_cuda/numba/cuda/core/targetconfig.py @@ -12,6 +12,7 @@ from types import MappingProxyType from numba.cuda import utils +from numba.cuda import HAS_NUMBA class Option: @@ -48,9 +49,9 @@ def doc(self): return self._doc -try: +if HAS_NUMBA: from numba.core.targetconfig import ConfigStack, _FlagsStack -except ImportError: +else: class _FlagsStack(utils.ThreadLocalStack, stack_name="flags"): pass diff --git a/numba_cuda/numba/cuda/core/tracing.py b/numba_cuda/numba/cuda/core/tracing.py index 8736535ed..ea8059fc9 100644 --- a/numba_cuda/numba/cuda/core/tracing.py +++ b/numba_cuda/numba/cuda/core/tracing.py @@ -8,7 +8,7 @@ from functools import wraps from itertools import chain -from numba.core import config +from numba.cuda.core import config class TLS(threading.local): diff --git a/numba_cuda/numba/cuda/core/transforms.py b/numba_cuda/numba/cuda/core/transforms.py index aa29f3ec4..90611557e 100644 --- a/numba_cuda/numba/cuda/core/transforms.py +++ b/numba_cuda/numba/cuda/core/transforms.py @@ -14,7 +14,7 @@ find_top_level_loops, ) from numba.cuda.core import ir -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.core import ir_utils from numba.cuda.core.analysis import compute_use_defs diff --git a/numba_cuda/numba/cuda/core/typed_passes.py b/numba_cuda/numba/cuda/core/typed_passes.py index 51b39f67c..c0d9f1455 100644 --- a/numba_cuda/numba/cuda/core/typed_passes.py +++ b/numba_cuda/numba/cuda/core/typed_passes.py @@ -400,7 +400,7 @@ def run_pass(self, state): @register_pass(mutates_CFG=True, analysis_only=False) class NativeLowering(BaseNativeLowering): """Lowering pass for a native function IR described solely in terms of - Numba's standard `numba.core.ir` nodes.""" + Numba's standard `numba.cuda.core.ir` nodes.""" _name = "native_lowering" diff --git a/numba_cuda/numba/cuda/core/unsafe/eh.py b/numba_cuda/numba/cuda/core/unsafe/eh.py index 851ec3337..0789d5812 100644 --- a/numba_cuda/numba/cuda/core/unsafe/eh.py +++ b/numba_cuda/numba/cuda/core/unsafe/eh.py @@ -6,7 +6,7 @@ """ from numba.cuda import types -from numba.core import errors +from numba.cuda.core import errors from numba.cuda import cgutils from numba.cuda.extending import intrinsic diff --git a/numba_cuda/numba/cuda/core/untyped_passes.py b/numba_cuda/numba/cuda/core/untyped_passes.py index 4006854ce..b0bd419e7 100644 --- a/numba_cuda/numba/cuda/core/untyped_passes.py +++ b/numba_cuda/numba/cuda/core/untyped_passes.py @@ -12,6 +12,7 @@ SSACompliantMixin, register_pass, ) +from numba import cuda from numba.cuda.core import postproc, bytecode, transforms, inline_closurecall from numba.cuda.core import ( errors, @@ -632,8 +633,6 @@ def __init__(self): FunctionPass.__init__(self) def run_pass(self, state): - from numba import njit - func_ir = state.func_ir mutated = False for idx, blk in func_ir.blocks.items(): @@ -669,7 +668,7 @@ def run_pass(self, state): continue pyfunc = convert_code_obj_to_function(node, func_ir) - func = njit()(pyfunc) + func = cuda.jit()(pyfunc) new_node = ir.Global( node.code.co_name, func, stmt.loc ) diff --git a/numba_cuda/numba/cuda/cpython/builtins.py b/numba_cuda/numba/cuda/cpython/builtins.py index 2c469a0e2..8a5eeac09 100644 --- a/numba_cuda/numba/cuda/cpython/builtins.py +++ b/numba_cuda/numba/cuda/cpython/builtins.py @@ -29,7 +29,7 @@ ) from numba.cuda.typing.templates import AbstractTemplate, signature from numba.cuda.typing.templates import infer_global -from numba.misc.special import literal_unroll +from numba.cuda.misc.special import literal_unroll from numba.cuda.typing.asnumbatype import as_numba_type from numba.cuda.typing.builtins import IndexValue, IndexValueType diff --git a/numba_cuda/numba/cuda/cpython/listobj.py b/numba_cuda/numba/cuda/cpython/listobj.py index de93b23ee..ee059a8e7 100644 --- a/numba_cuda/numba/cuda/cpython/listobj.py +++ b/numba_cuda/numba/cuda/cpython/listobj.py @@ -9,7 +9,7 @@ from llvmlite import ir from numba.cuda import types -from numba.core import errors +from numba.cuda.core import errors from numba.cuda import cgutils from numba.cuda.core.imputils import ( Registry, @@ -19,9 +19,8 @@ RefType, ) from numba.cuda.extending import overload_method, overload -from numba.misc import quicksort from numba.cuda.cpython import slicing -from numba import literal_unroll +from numba.cuda.misc.special import literal_unroll registry = Registry("listobj") lower = registry.lower @@ -1132,17 +1131,6 @@ def gt(a, b): return a > b -sort_forwards = quicksort.make_jit_quicksort().run_quicksort -sort_backwards = quicksort.make_jit_quicksort(lt=gt).run_quicksort - -arg_sort_forwards = quicksort.make_jit_quicksort( - is_argsort=True, is_list=True -).run_quicksort -arg_sort_backwards = quicksort.make_jit_quicksort( - is_argsort=True, lt=gt, is_list=True -).run_quicksort - - def _sort_check_reverse(reverse): if isinstance(reverse, types.Omitted): rty = reverse.value @@ -1168,35 +1156,6 @@ def _sort_check_key(key): raise errors.TypingError(msg) -@overload_method(types.List, "sort") -def ol_list_sort(lst, key=None, reverse=False): - _sort_check_key(key) - _sort_check_reverse(reverse) - - if cgutils.is_nonelike(key): - KEY = False - sort_f = sort_forwards - sort_b = sort_backwards - elif isinstance(key, types.Dispatcher): - KEY = True - sort_f = arg_sort_forwards - sort_b = arg_sort_backwards - - def impl(lst, key=None, reverse=False): - if KEY is True: - _lst = [key(x) for x in lst] - else: - _lst = lst - if reverse is False or reverse == 0: - tmp = sort_f(_lst) - else: - tmp = sort_b(_lst) - if KEY is True: - lst[:] = [lst[i] for i in tmp] - - return impl - - @overload(sorted) def ol_sorted(iterable, key=None, reverse=False): if not isinstance(iterable, types.IterableType): diff --git a/numba_cuda/numba/cuda/cpython/numbers.py b/numba_cuda/numba/cuda/cpython/numbers.py index e5dd6e224..d3a24b934 100644 --- a/numba_cuda/numba/cuda/cpython/numbers.py +++ b/numba_cuda/numba/cuda/cpython/numbers.py @@ -12,7 +12,7 @@ from numba.cuda.core.imputils import impl_ret_untracked, Registry from numba.cuda import types -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.extending import overload_method from numba.cuda.cpython.unsafe.numbers import viewer from numba.cuda import cgutils, typing diff --git a/numba_cuda/numba/cuda/cpython/unsafe/numbers.py b/numba_cuda/numba/cuda/cpython/unsafe/numbers.py index 4cff57200..82c878a81 100644 --- a/numba_cuda/numba/cuda/cpython/unsafe/numbers.py +++ b/numba_cuda/numba/cuda/cpython/unsafe/numbers.py @@ -4,7 +4,7 @@ """This module provides the unsafe things for targets/numbers.py""" from numba.cuda import types -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.extending import intrinsic from llvmlite import ir diff --git a/numba_cuda/numba/cuda/cpython/unsafe/tuple.py b/numba_cuda/numba/cuda/cpython/unsafe/tuple.py index 50f1d709e..e3e47c981 100644 --- a/numba_cuda/numba/cuda/cpython/unsafe/tuple.py +++ b/numba_cuda/numba/cuda/cpython/unsafe/tuple.py @@ -7,7 +7,7 @@ """ from numba.cuda import types, typing -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.cgutils import alloca_once from numba.cuda.extending import intrinsic diff --git a/numba_cuda/numba/cuda/extending.py b/numba_cuda/numba/cuda/extending.py index a9642917f..11de8d096 100644 --- a/numba_cuda/numba/cuda/extending.py +++ b/numba_cuda/numba/cuda/extending.py @@ -10,7 +10,7 @@ import collections import functools -from numba.core import errors +from numba.cuda.core import errors from numba.cuda import types, utils, config # # Exported symbols diff --git a/numba_cuda/numba/cuda/lowering.py b/numba_cuda/numba/cuda/lowering.py index 16a557dbc..0f47ea6b9 100644 --- a/numba_cuda/numba/cuda/lowering.py +++ b/numba_cuda/numba/cuda/lowering.py @@ -1,7 +1,6 @@ # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: BSD-2-Clause -import numba from collections import namedtuple, defaultdict import operator import warnings @@ -9,6 +8,7 @@ from llvmlite import ir as llvm_ir +from numba.cuda import HAS_NUMBA from numba.cuda.core import ir from numba.cuda import debuginfo, cgutils, utils, typing, types from numba.cuda.core import ( @@ -32,12 +32,7 @@ from numba.cuda.core.environment import Environment from numba.cuda.core.analysis import compute_use_defs, must_use_alloca from numba.cuda.misc.firstlinefinder import get_func_body_first_lineno -from numba import version_info - -numba_version = version_info.short -del version_info -if numba_version > (0, 60): - from numba.cuda.misc.coverage_support import get_registered_loc_notify +from numba.cuda.misc.coverage_support import get_registered_loc_notify _VarArgItem = namedtuple("_VarArgItem", ("vararg", "index")) @@ -96,9 +91,8 @@ def __init__(self, context, library, fndesc, func_ir, metadata=None): directives_only=directives_only, ) - if numba_version > (0, 60): - # Loc notify objects - self._loc_notify_registry = get_registered_loc_notify() + # Loc notify objects + self._loc_notify_registry = get_registered_loc_notify() # Subclass initialization self.init() @@ -174,9 +168,8 @@ def post_lower(self): Called after all blocks are lowered """ self.debuginfo.finalize() - if numba_version > (0, 60): - for notify in self._loc_notify_registry: - notify.close() + for notify in self._loc_notify_registry: + notify.close() def pre_block(self, block): """ @@ -369,11 +362,8 @@ def notify_loc(self, loc: ir.Loc) -> None: """Called when a new instruction with the given `loc` is about to be lowered. """ - if numba_version > (0, 60): - for notify_obj in self._loc_notify_registry: - notify_obj.notify(loc) - else: - pass + for notify_obj in self._loc_notify_registry: + notify_obj.notify(loc) def debug_print(self, msg): if config.DEBUG_JIT: @@ -1880,7 +1870,14 @@ def _lit_or_omitted(value): """Returns a Literal instance if the type of value is supported; otherwise, return `Omitted(value)`. """ + typing_errors = LiteralTypingError + if HAS_NUMBA: + from numba.core.errors import ( + LiteralTypingError as CoreLiteralTypingError, + ) + + typing_errors = (LiteralTypingError, CoreLiteralTypingError) try: return types.literal(value) - except (LiteralTypingError, numba.core.errors.LiteralTypingError): + except typing_errors: return types.Omitted(value) diff --git a/numba_cuda/numba/cuda/memory_management/nrt_context.py b/numba_cuda/numba/cuda/memory_management/nrt_context.py index dd8bd3c66..208412d3b 100644 --- a/numba_cuda/numba/cuda/memory_management/nrt_context.py +++ b/numba_cuda/numba/cuda/memory_management/nrt_context.py @@ -6,7 +6,7 @@ from llvmlite import ir from numba.cuda import types -from numba.core import errors +from numba.cuda.core import errors from numba.cuda import cgutils, config from numba.cuda.utils import PYVERSION diff --git a/numba_cuda/numba/cuda/misc/coverage_support.py b/numba_cuda/numba/cuda/misc/coverage_support.py index 1167fd696..b3fa17a52 100644 --- a/numba_cuda/numba/cuda/misc/coverage_support.py +++ b/numba_cuda/numba/cuda/misc/coverage_support.py @@ -21,7 +21,7 @@ def get_registered_loc_notify() -> Sequence["NotifyLocBase"]: """ Returns a list of the registered NotifyLocBase instances. """ - if not config.JIT_COVERAGE: + if hasattr(config, "JIT_COVERAGE") and not config.JIT_COVERAGE: # Coverage disabled. return [] return list( @@ -32,7 +32,7 @@ def get_registered_loc_notify() -> Sequence["NotifyLocBase"]: class NotifyLocBase(ABC): - """Interface for notifying visiting of a ``numba.core.ir.Loc``.""" + """Interface for notifying visiting of a ``numba.cuda.core.ir.Loc``.""" @abstractmethod def notify(self, loc: ir.Loc) -> None: diff --git a/numba_cuda/numba/cuda/misc/gdb_hook.py b/numba_cuda/numba/cuda/misc/gdb_hook.py index ee945a3b0..2c8bd3c8d 100644 --- a/numba_cuda/numba/cuda/misc/gdb_hook.py +++ b/numba_cuda/numba/cuda/misc/gdb_hook.py @@ -7,7 +7,7 @@ from llvmlite import ir from numba.cuda import types, config -from numba.core import errors +from numba.cuda.core import errors from numba.cuda import cgutils, utils from numba.cuda.misc.special import gdb, gdb_init, gdb_breakpoint from numba.cuda.extending import overload, intrinsic diff --git a/numba_cuda/numba/cuda/np/arraymath.py b/numba_cuda/numba/cuda/np/arraymath.py index cbcef34db..7fcf2458e 100644 --- a/numba_cuda/numba/cuda/np/arraymath.py +++ b/numba_cuda/numba/cuda/np/arraymath.py @@ -5111,75 +5111,6 @@ def np_setdiff1d_impl(ar1, ar2, assume_unique=False): return np_setdiff1d_impl -@overload(np.in1d) -def jit_np_in1d(ar1, ar2, assume_unique=False, invert=False): - if not (type_can_asarray(ar1) or type_can_asarray(ar2)): - raise TypingError("in1d: first two args must be array-like") - if not isinstance(assume_unique, (types.Boolean, bool)): - raise TypingError('in1d: Argument "assume_unique" must be boolean') - if not isinstance(invert, (types.Boolean, bool)): - raise TypingError('in1d: Argument "invert" must be boolean') - - def np_in1d_impl(ar1, ar2, assume_unique=False, invert=False): - # https://github.com/numpy/numpy/blob/03b62604eead0f7d279a5a4c094743eb29647368/numpy/lib/arraysetops.py#L525 # noqa: E501 - - # Ravel both arrays, behavior for the first array could be different - ar1 = np.asarray(ar1).ravel() - ar2 = np.asarray(ar2).ravel() - - # This code is run when it would make the code significantly faster - # Sorting is also not guaranteed to work on objects but numba does - # not support object arrays. - if len(ar2) < 10 * len(ar1) ** 0.145: - if invert: - mask = np.ones(len(ar1), dtype=np.bool_) - for a in ar2: - mask &= ar1 != a - else: - mask = np.zeros(len(ar1), dtype=np.bool_) - for a in ar2: - mask |= ar1 == a - return mask - - # Otherwise use sorting - if not assume_unique: - # Equivalent to ar1, inv_idx = np.unique(ar1, return_inverse=True) - # https://github.com/numpy/numpy/blob/03b62604eead0f7d279a5a4c094743eb29647368/numpy/lib/arraysetops.py#L358C8-L358C8 # noqa: E501 - order1 = np.argsort(ar1) - aux = ar1[order1] - mask = np.empty(aux.shape, dtype=np.bool_) - mask[:1] = True - mask[1:] = aux[1:] != aux[:-1] - ar1 = aux[mask] - imask = np.cumsum(mask) - 1 - inv_idx = np.empty(mask.shape, dtype=np.intp) - inv_idx[order1] = imask - ar2 = np.unique(ar2) - - ar = np.concatenate((ar1, ar2)) - # We need this to be a stable sort, so always use 'mergesort' - # here. The values from the first array should always come before - # the values from the second array. - order = ar.argsort(kind="mergesort") - sar = ar[order] - flag = np.empty(sar.size, np.bool_) - if invert: - flag[:-1] = sar[1:] != sar[:-1] - else: - flag[:-1] = sar[1:] == sar[:-1] - flag[-1:] = invert - ret = np.empty(ar.shape, dtype=np.bool_) - ret[order] = flag - - # return ret[:len(ar1)] - if assume_unique: - return ret[: len(ar1)] - else: - return ret[inv_idx] - - return np_in1d_impl - - @overload(np.isin) def jit_np_isin(element, test_elements, assume_unique=False, invert=False): if not (type_can_asarray(element) or type_can_asarray(test_elements)): diff --git a/numba_cuda/numba/cuda/np/arrayobj.py b/numba_cuda/numba/cuda/np/arrayobj.py index fcf979f4f..de8834e91 100644 --- a/numba_cuda/numba/cuda/np/arrayobj.py +++ b/numba_cuda/numba/cuda/np/arrayobj.py @@ -16,9 +16,9 @@ import numpy as np -from numba import pndindex, literal_unroll +from numba.cuda.misc.special import literal_unroll from numba.cuda import types, typing -from numba.core import errors +from numba.cuda.core import errors from numba.cuda import cgutils, extending from numba.cuda.np.numpy_support import ( as_dtype, @@ -54,7 +54,6 @@ intrinsic, overload_attribute, ) -from numba.misc import quicksort, mergesort from numba.cuda.cpython import slicing from numba.cuda.cpython.unsafe.tuple import ( tuple_setitem, @@ -4529,7 +4528,6 @@ def iternext_numpy_nditer(context, builder, sig, args, result): nditer.iternext_specific(context, builder, arrty, arr, result) -@lower(pndindex, types.VarArg(types.Integer)) @lower(np.ndindex, types.VarArg(types.Integer)) def make_array_ndindex(context, builder, sig, args): """ndindex(*shape)""" @@ -4546,7 +4544,6 @@ def make_array_ndindex(context, builder, sig, args): return impl_ret_borrowed(context, builder, sig.return_type, res) -@lower(pndindex, types.BaseTuple) @lower(np.ndindex, types.BaseTuple) def make_array_ndindex_tuple(context, builder, sig, args): """ndindex(shape)""" @@ -7127,18 +7124,7 @@ def get_sort_func(kind, lt_impl, is_argsort=False): try: return _sorts[key] except KeyError: - if kind == "quicksort": - sort = quicksort.make_jit_quicksort( - lt=lt_impl, is_argsort=is_argsort, is_np_array=True - ) - func = sort.run_quicksort - elif kind == "mergesort": - sort = mergesort.make_jit_mergesort( - lt=lt_impl, is_argsort=is_argsort - ) - func = sort.run_mergesort - _sorts[key] = func - return func + raise errors.NumbaError("Sort kind %s not supported" % kind) def lt_implementation(dtype): @@ -7150,21 +7136,6 @@ def lt_implementation(dtype): return default_lt -@lower("array.sort", types.Array) -def array_sort(context, builder, sig, args): - arytype = sig.args[0] - - sort_func = get_sort_func( - kind="quicksort", lt_impl=lt_implementation(arytype.dtype) - ) - - def array_sort_impl(arr): - # Note we clobber the return value - sort_func(arr) - - return context.compile_internal(builder, array_sort_impl, sig, args) - - @overload(np.sort) def impl_np_sort(a): if not type_can_asarray(a): diff --git a/numba_cuda/numba/cuda/np/math/numbers.py b/numba_cuda/numba/cuda/np/math/numbers.py index 8dd3a92b0..cf0875b9c 100644 --- a/numba_cuda/numba/cuda/np/math/numbers.py +++ b/numba_cuda/numba/cuda/np/math/numbers.py @@ -11,7 +11,7 @@ from numba.cuda.core.imputils import impl_ret_untracked from numba.cuda import typing, types, cgutils -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.cpython.unsafe.numbers import viewer diff --git a/numba_cuda/numba/cuda/np/npyfuncs.py b/numba_cuda/numba/cuda/np/npyfuncs.py index 5164d861f..806dd301d 100644 --- a/numba_cuda/numba/cuda/np/npyfuncs.py +++ b/numba_cuda/numba/cuda/np/npyfuncs.py @@ -15,7 +15,7 @@ from numba.cuda.extending import overload from numba.cuda.core.imputils import impl_ret_untracked from numba.cuda import types -from numba.core import errors +from numba.cuda.core import errors from numba.cuda import cgutils, typing from numba.cuda.np import npdatetime from numba.cuda.extending import register_jitable diff --git a/numba_cuda/numba/cuda/np/npyimpl.py b/numba_cuda/numba/cuda/np/npyimpl.py index fb42cee8c..472f55645 100644 --- a/numba_cuda/numba/cuda/np/npyimpl.py +++ b/numba_cuda/numba/cuda/np/npyimpl.py @@ -34,7 +34,7 @@ from numba.cuda.typing import npydecl from numba.cuda.extending import overload, intrinsic -from numba.core import errors +from numba.cuda.core import errors registry = Registry("npyimpl") diff --git a/numba_cuda/numba/cuda/np/numpy_support.py b/numba_cuda/numba/cuda/np/numpy_support.py index f850b2bc8..69d048efa 100644 --- a/numba_cuda/numba/cuda/np/numpy_support.py +++ b/numba_cuda/numba/cuda/np/numpy_support.py @@ -7,7 +7,7 @@ import numpy as np from numba.cuda import types -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.typing.templates import signature from numba.cuda.np import npdatetime_helpers from numba.cuda.core.errors import TypingError diff --git a/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py b/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py index 4cd48baf0..5d974f8da 100644 --- a/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py +++ b/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py @@ -9,9 +9,9 @@ from numpy.polynomial import polynomial as poly from numpy.polynomial import polyutils as pu -from numba import literal_unroll +from numba.cuda.misc.special import literal_unroll from numba.cuda import types -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.extending import overload from numba.cuda.np.numpy_support import type_can_asarray, as_dtype, from_dtype diff --git a/numba_cuda/numba/cuda/np/ufunc/decorators.py b/numba_cuda/numba/cuda/np/ufunc/decorators.py index 80cb11c6a..f71cecbf5 100644 --- a/numba_cuda/numba/cuda/np/ufunc/decorators.py +++ b/numba_cuda/numba/cuda/np/ufunc/decorators.py @@ -1,7 +1,7 @@ # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: BSD-2-Clause -from numba.core.registry import DelayedRegistry +from numba.cuda.core.registry import DelayedRegistry from numba.cuda.vectorizers import CUDAVectorize, CUDAGUFuncVectorize diff --git a/numba_cuda/numba/cuda/random.py b/numba_cuda/numba/cuda/random.py index 47a76b248..e8e620d5c 100644 --- a/numba_cuda/numba/cuda/random.py +++ b/numba_cuda/numba/cuda/random.py @@ -3,18 +3,21 @@ import math -from numba import ( - cuda, +from numba import cuda +from numba.cuda import ( float32, float64, uint32, int64, uint64, - from_dtype, - jit, + HAS_NUMBA, ) +from numba.cuda.np.numpy_support import from_dtype from numba.cuda import config +if HAS_NUMBA: + from numba import jit + import numpy as np # This implementation is based upon the xoroshiro128+ and splitmix64 algorithms diff --git a/numba_cuda/numba/cuda/target.py b/numba_cuda/numba/cuda/target.py index 6faf15571..3425bd181 100644 --- a/numba_cuda/numba/cuda/target.py +++ b/numba_cuda/numba/cuda/target.py @@ -9,7 +9,7 @@ import importlib.util from numba.cuda import types - +from numba.cuda import HAS_NUMBA from numba.core.compiler_lock import global_compiler_lock from numba.cuda.core.errors import NumbaWarning from numba.cuda.core.base import BaseContext @@ -64,9 +64,7 @@ def resolve_value_type(self, val): from numba.cuda.dispatcher import CUDADispatcher from numba.core.dispatcher import Dispatcher - try: - from numba.core.dispatcher import Dispatcher - + if HAS_NUMBA: if isinstance(val, Dispatcher) and not isinstance( val, CUDADispatcher ): @@ -88,8 +86,6 @@ def resolve_value_type(self, val): # duplicated copy of the same function. val.__dispatcher = disp val = disp - except ImportError: - pass # continue with parent logic return super(CUDATypingContext, self).resolve_value_type(val) diff --git a/numba_cuda/numba/cuda/testing.py b/numba_cuda/numba/cuda/testing.py index 196d08897..733cd96fb 100644 --- a/numba_cuda/numba/cuda/testing.py +++ b/numba_cuda/numba/cuda/testing.py @@ -18,6 +18,7 @@ from io import StringIO import unittest import numpy as np +from numba.cuda import HAS_NUMBA if PYVERSION >= (3, 10): from filecheck.matcher import Matcher @@ -189,6 +190,11 @@ def skip_on_cudasim(reason): return unittest.skipIf(config.ENABLE_CUDASIM, reason) +def skip_on_standalone_numba_cuda(reason): + """Skip this test if running on standalone numba_cuda""" + return unittest.skipIf(not HAS_NUMBA, reason) + + def skip_unless_cudasim(reason): """Skip this test if running on CUDA hardware""" return unittest.skipUnless(config.ENABLE_CUDASIM, reason) diff --git a/numba_cuda/numba/cuda/tests/core/serialize_usecases.py b/numba_cuda/numba/cuda/tests/core/serialize_usecases.py index bd0fcfc51..37e676b6e 100644 --- a/numba_cuda/numba/cuda/tests/core/serialize_usecases.py +++ b/numba_cuda/numba/cuda/tests/core/serialize_usecases.py @@ -11,6 +11,7 @@ import numpy as np import numpy.random as nprand +# This does not need a guard, it's already guarded at the import site from numba import jit diff --git a/numba_cuda/numba/cuda/tests/core/test_serialize.py b/numba_cuda/numba/cuda/tests/core/test_serialize.py index f5db16869..d2da82c47 100644 --- a/numba_cuda/numba/cuda/tests/core/test_serialize.py +++ b/numba_cuda/numba/cuda/tests/core/test_serialize.py @@ -10,21 +10,20 @@ import unittest from multiprocessing import get_context +from numba.cuda import HAS_NUMBA import numba -from numba.core.errors import TypingError -from numba.cuda.tests.support import TestCase -from numba.cuda.cloudpickle import dumps, loads -try: +if HAS_NUMBA: + from numba.core.errors import TypingError from numba.core.target_extension import resolve_dispatcher_from_str -except ImportError: - resolve_dispatcher_from_str = None +else: + from numba.cuda.core.errors import TypingError +from numba.cuda.tests.support import TestCase +from numba.cuda.cloudpickle import dumps, loads +from numba.cuda.testing import skip_on_standalone_numba_cuda -@unittest.skipIf( - resolve_dispatcher_from_str is None, - "numba.core.target_extension not available", -) +@skip_on_standalone_numba_cuda class TestDispatcherPickling(TestCase): def run_with_protocols(self, meth, *args, **kwargs): for proto in range(pickle.HIGHEST_PROTOCOL + 1): @@ -305,6 +304,8 @@ class Klass: self.assertEqual(proc.exitcode, 0) def test_dynamic_class_issue_7356(self): + import numba + cfunc = numba.njit(issue_7356) self.assertEqual(cfunc(), (100, 100)) diff --git a/numba_cuda/numba/cuda/tests/cudapy/cache_with_cpu_usecases.py b/numba_cuda/numba/cuda/tests/cudapy/cache_with_cpu_usecases.py index 8c9cce605..ce1d9a041 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/cache_with_cpu_usecases.py +++ b/numba_cuda/numba/cuda/tests/cudapy/cache_with_cpu_usecases.py @@ -3,8 +3,9 @@ import sys -from numba import cuda, njit -from numba.cuda.testing import CUDATestCase +from numba import cuda +from numba.cuda import HAS_NUMBA +from numba.cuda.testing import CUDATestCase, skip_on_standalone_numba_cuda from numba.cuda.tests.cudapy.cache_usecases import CUDAUseCase, UseCase @@ -22,10 +23,14 @@ def target_shared_assign(r, x): assign_cuda_kernel = cuda.jit(cache=True)(target_shared_assign) assign_cuda = CUDAUseCase(assign_cuda_kernel) -assign_cpu_jitted = njit(cache=True)(target_shared_assign) -assign_cpu = CPUUseCase(assign_cpu_jitted) +if HAS_NUMBA: + from numba import njit + assign_cpu_jitted = njit(cache=True)(target_shared_assign) + assign_cpu = CPUUseCase(assign_cpu_jitted) + +@skip_on_standalone_numba_cuda class _TestModule(CUDATestCase): """ Tests for functionality of this module's functions. diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_array_alignment.py b/numba_cuda/numba/cuda/tests/cudapy/test_array_alignment.py index 662796be4..24d14dd58 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_array_alignment.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_array_alignment.py @@ -5,7 +5,12 @@ import itertools import numpy as np from numba import cuda -from numba.core.errors import TypingError +from numba.cuda import HAS_NUMBA + +if HAS_NUMBA: + from numba.core.errors import TypingError +else: + from numba.cuda.core.errors import TypingError from numba.cuda.testing import ( CUDATestCase, skip_on_cudasim, diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_bfloat16_bindings.py b/numba_cuda/numba/cuda/tests/cudapy/test_bfloat16_bindings.py index 4916059a8..de3d44c50 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_bfloat16_bindings.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_bfloat16_bindings.py @@ -4,7 +4,7 @@ from collections import OrderedDict import bisect -import numba.cuda as cuda +from numba import cuda from numba.cuda.testing import unittest, CUDATestCase import numpy as np import operator diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_caching.py b/numba_cuda/numba/cuda/tests/cudapy/test_caching.py index 9785cb54f..1bfd03595 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_caching.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_caching.py @@ -18,6 +18,7 @@ skip_unless_cc_60, skip_if_cudadevrt_missing, test_data_dir, + skip_on_standalone_numba_cuda, ) from numba.cuda.tests.support import ( TestCase, @@ -415,6 +416,7 @@ def tearDown(self): CUDATestCase.tearDown(self) DispatcherCacheUsecasesTest.tearDown(self) + @skip_on_standalone_numba_cuda def test_cpu_and_cuda_targets(self): # The same function jitted for CPU and CUDA targets should maintain # separate caches for each target. @@ -440,6 +442,7 @@ def test_cpu_and_cuda_targets(self): self.check_hits(f_cpu.func, 0, 2) self.check_hits(f_cuda.func, 0, 2) + @skip_on_standalone_numba_cuda def test_cpu_and_cuda_reuse(self): # Existing cache files for the CPU and CUDA targets are reused. mod = self.import_module() diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_cuda_array_interface.py b/numba_cuda/numba/cuda/tests/cudapy/test_cuda_array_interface.py index 15a7b4419..d9a642144 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_cuda_array_interface.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_cuda_array_interface.py @@ -3,7 +3,7 @@ import numpy as np -from numba import vectorize, guvectorize +from numba.cuda import vectorize, guvectorize from numba import cuda from numba.cuda.testing import unittest, CUDATestCase, ForeignArray from numba.cuda.testing import skip_on_cudasim, skip_if_external_memmgr diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo_types.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo_types.py index e6e48f7b4..916563335 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo_types.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo_types.py @@ -1,7 +1,7 @@ # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: BSD-2-Clause -import numba.cuda as cuda +from numba import cuda from numba.cuda.testing import CUDATestCase, skip_on_cudasim import llvmlite from numba.cuda import types diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_device_func.py b/numba_cuda/numba/cuda/tests/cudapy/test_device_func.py index f941a020c..3fa62728b 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_device_func.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_device_func.py @@ -13,11 +13,16 @@ unittest, CUDATestCase, ) -from numba import cuda, jit +from numba import cuda from numba.cuda import float32, int32, types from numba.cuda.core.errors import TypingError from numba.cuda.tests.support import skip_unless_cffi +from numba.cuda.testing import skip_on_standalone_numba_cuda from types import ModuleType +from numba.cuda import HAS_NUMBA + +if HAS_NUMBA: + from numba import jit class TestDeviceFunc(CUDATestCase): @@ -72,6 +77,7 @@ def add_kernel(ary): add_kernel[1, ary.size](ary) np.testing.assert_equal(expect, ary) + @skip_on_standalone_numba_cuda def test_cpu_dispatcher(self): # Test correct usage @jit @@ -81,6 +87,7 @@ def add(a, b): self._check_cpu_dispatcher(add) @skip_on_cudasim("not supported in cudasim") + @skip_on_standalone_numba_cuda def test_cpu_dispatcher_invalid(self): # Test invalid usage # Explicit signature disables compilation, which also disable @@ -96,6 +103,7 @@ def add(a, b): expected = re.compile(msg) self.assertTrue(expected.search(str(raises.exception)) is not None) + @skip_on_standalone_numba_cuda def test_cpu_dispatcher_other_module(self): @jit def add(a, b): diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_enums.py b/numba_cuda/numba/cuda/tests/cudapy/test_enums.py index 12c5161b4..d7f05a7f7 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_enums.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_enums.py @@ -8,9 +8,18 @@ import numpy as np from numba.cuda import int16, int32 -from numba import cuda, vectorize, njit +from numba.cuda import vectorize, HAS_NUMBA + +if HAS_NUMBA: + from numba import njit +from numba import cuda from numba.cuda import types -from numba.cuda.testing import unittest, CUDATestCase, skip_on_cudasim +from numba.cuda.testing import ( + unittest, + CUDATestCase, + skip_on_cudasim, + skip_on_standalone_numba_cuda, +) from numba.cuda.tests.enum_usecases import ( Color, Shape, @@ -57,6 +66,7 @@ def f(out): f(expected) self.assertPreciseEqual(expected, got) + @skip_on_standalone_numba_cuda def test_return_from_device_func(self): @njit def inner(pred): diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_extending.py b/numba_cuda/numba/cuda/tests/cudapy/test_extending.py index f0e920b83..df5afccb1 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_extending.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_extending.py @@ -7,7 +7,9 @@ import numpy as np import os -from numba import cuda, njit +from numba import cuda +from numba.cuda import HAS_NUMBA +from numba.cuda.testing import skip_on_standalone_numba_cuda from numba.cuda import types from numba.cuda import config from numba.cuda.extending import overload @@ -30,6 +32,12 @@ def width(self): return self.hi - self.lo +if HAS_NUMBA: + from numba import njit +else: + njit = None + + @njit def interval_width(interval): return interval.width @@ -138,6 +146,7 @@ def f(r, x): np.testing.assert_allclose(r[0], x[1] - x[0]) + @skip_on_standalone_numba_cuda def test_extension_type_as_arg(self): @cuda.jit def f(r, x): @@ -151,6 +160,7 @@ def f(r, x): np.testing.assert_allclose(r[0], x[1] - x[0]) + @skip_on_standalone_numba_cuda def test_extension_type_as_retvalue(self): @cuda.jit def f(r, x): diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_gufunc.py b/numba_cuda/numba/cuda/tests/cudapy/test_gufunc.py index 0d7b1c075..223b9ad2c 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_gufunc.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_gufunc.py @@ -5,7 +5,7 @@ from collections import namedtuple from numba.cuda import void, int32, float32, float64 -from numba import guvectorize +from numba.cuda import guvectorize from numba import cuda from numba.cuda.testing import skip_on_cudasim, CUDATestCase import unittest diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_gufunc_scalar.py b/numba_cuda/numba/cuda/tests/cudapy/test_gufunc_scalar.py index 819be81b8..64769f585 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_gufunc_scalar.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_gufunc_scalar.py @@ -8,7 +8,7 @@ """ import numpy as np -from numba import guvectorize, cuda +from numba import cuda, guvectorize from numba.cuda.testing import skip_on_cudasim, CUDATestCase import unittest diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_intrinsics.py b/numba_cuda/numba/cuda/tests/cudapy/test_intrinsics.py index e01f87a22..93fbec026 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_intrinsics.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_intrinsics.py @@ -7,7 +7,12 @@ import re from numba import cuda from numba.cuda import int64 -from numba.core.errors import TypingError +from numba.cuda import HAS_NUMBA + +if HAS_NUMBA: + from numba.core.errors import TypingError +else: + from numba.cuda.core.errors import TypingError from numba.cuda.types import f2 from numba.cuda.testing import ( unittest, diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_ir.py b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py index 7fac682cb..634ac15ec 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_ir.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py @@ -2,14 +2,18 @@ # SPDX-License-Identifier: BSD-2-Clause import unittest -from numba.cuda.testing import CUDATestCase +from numba.cuda.testing import CUDATestCase, skip_on_cudasim import warnings import numpy as np -from numba import objmode +from numba.cuda import HAS_NUMBA + +if HAS_NUMBA: + from numba import objmode + from numba import njit from numba.cuda.core import ir from numba.cuda import compiler -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.core.compiler import ( CompilerBase, ) @@ -23,7 +27,7 @@ IRProcessing, ReconstructSSA, ) -from numba import njit +from numba.cuda.testing import skip_on_standalone_numba_cuda def requires_fully_vendored_ir(fn): @@ -374,7 +378,7 @@ def gen_block(): self.check(a, same=[b], different=[c]) - @requires_fully_vendored_ir + @skip_on_cudasim def test_functionir(self): def run_frontend(x): return compiler.run_frontend(x, emit_dels=True) @@ -396,8 +400,9 @@ def foo(a, b, c=12, d=1j, e=None): if np.abs(i) > 0: k = h / i l = np.arange(1, c + 1) - with objmode(): - print(e, k) + if HAS_NUMBA: + with objmode(): + print(e, k) m = np.sqrt(l - g) if np.abs(m[0]) < 1: n = 0 @@ -411,9 +416,10 @@ def foo(a, b, c=12, d=1j, e=None): for r in range(len(p)): q.append(p[r]) if r > 4 + 1: - with objmode(s="intp", t="complex128"): - s = 123 - t = 5 + if HAS_NUMBA: + with objmode(s="intp", t="complex128"): + s = 123 + t = 5 if s > 122: t += s t += q[0] + _GLOBAL @@ -514,8 +520,8 @@ def baz(a, b): check_diffstr(tmp, ["c + b", "b + c"]) -@requires_fully_vendored_ir class TestIRPedanticChecks(CUDATestCase): + @skip_on_standalone_numba_cuda def test_var_in_scope_assumption(self): # Create a pass that clears ir.Scope in ir.Block @register_pass(mutates_CFG=False, analysis_only=False) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_localmem.py b/numba_cuda/numba/cuda/tests/cudapy/test_localmem.py index 0099a532d..8f264be52 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_localmem.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_localmem.py @@ -6,7 +6,12 @@ from numba import cuda from numba.cuda import int32, complex128, void from numba.cuda import types -from numba.core.errors import TypingError +from numba.cuda import HAS_NUMBA + +if HAS_NUMBA: + from numba.core.errors import TypingError +else: + from numba.cuda.core.errors import TypingError from numba.cuda.testing import unittest, CUDATestCase, skip_on_cudasim from .extensions_usecases import test_struct_model_type, TestStruct diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_overload.py b/numba_cuda/numba/cuda/tests/cudapy/test_overload.py index ff8190671..2bba8f4e7 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_overload.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_overload.py @@ -1,12 +1,25 @@ # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: BSD-2-Clause -from numba import cuda, njit, types, version_info -from numba.core.errors import TypingError +from numba import cuda +from numba.cuda import types +from numba.cuda import HAS_NUMBA + +if HAS_NUMBA: + from numba.core.errors import TypingError + from numba import njit + import numba +else: + from numba.cuda.core.errors import TypingError from numba.cuda.extending import overload, overload_attribute from numba.cuda.typing.typeof import typeof from numba.core.typing.typeof import typeof as cpu_typeof -from numba.cuda.testing import CUDATestCase, skip_on_cudasim, unittest +from numba.cuda.testing import ( + CUDATestCase, + skip_on_cudasim, + unittest, + skip_on_standalone_numba_cuda, +) import numpy as np @@ -224,6 +237,7 @@ def check_overload(self, kernel, expected): cuda.jit(kernel)[1, 1](x) self.assertEqual(x[0], expected) + @skip_on_standalone_numba_cuda def check_overload_cpu(self, kernel, expected): x = np.ones(1, dtype=np.int32) njit(kernel)(x) @@ -324,10 +338,16 @@ def kernel(x): expected = CUDA_TARGET_OL_CALLS_TARGET_OL * CUDA_TARGET_OL self.check_overload(kernel, expected) - # Also check that the CPU overloads are used on the CPU + @skip_on_standalone_numba_cuda + def test_target_overloaded_calls_target_overloaded_cpu(self): + def kernel(x): + target_overloaded_calls_target_overloaded(x) + + # Check that the CPU overloads are used on the CPU expected = GENERIC_TARGET_OL_CALLS_TARGET_OL * GENERIC_TARGET_OL self.check_overload_cpu(kernel, expected) + @skip_on_standalone_numba_cuda def test_overload_attribute_target(self): MyDummy, MyDummyType = self.make_dummy_type() mydummy_type_cpu = cpu_typeof(MyDummy()) # For @njit (cpu) @@ -346,7 +366,7 @@ def imp(obj): # A different error is produced prior to version 0.60 # (the fixes in #9454 improved the message) # https://github.com/numba/numba/pull/9454 - if version_info[:2] < (0, 60): + if HAS_NUMBA and numba.version_info[:2] < (0, 60): msg = 'resolving type of attribute "cuda_only" of "x"' else: msg = "Unknown attribute 'cuda_only'" diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_random.py b/numba_cuda/numba/cuda/tests/cudapy/test_random.py index 99e7fc662..c99e29aa5 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_random.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_random.py @@ -7,7 +7,11 @@ from numba import cuda from numba.cuda.testing import unittest -from numba.cuda.testing import skip_on_cudasim, CUDATestCase +from numba.cuda.testing import ( + skip_on_cudasim, + CUDATestCase, + skip_on_standalone_numba_cuda, +) from numba.cuda.random import ( xoroshiro128p_uniform_float32, @@ -48,6 +52,7 @@ def rng_kernel_float64(states, out, count, distribution): out[idx] = xoroshiro128p_normal_float64(states, thread_id) +@skip_on_standalone_numba_cuda class TestCudaRandomXoroshiro128p(CUDATestCase): def test_create(self): states = cuda.random.create_xoroshiro128p_states(10, seed=1) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_sm.py b/numba_cuda/numba/cuda/tests/cudapy/test_sm.py index 85424fcdd..659a13cc7 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_sm.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_sm.py @@ -3,8 +3,11 @@ from numba import cuda from numba.cuda import int32, float64, void +from numba.cuda import HAS_NUMBA + +if HAS_NUMBA: + from numba.core.errors import TypingError as NumbaTypingError from numba.cuda.core.errors import TypingError -import numba.core.errors as numbaError from numba.cuda import types from numba.cuda.testing import unittest, CUDATestCase, skip_on_cudasim @@ -410,8 +413,12 @@ def unsupported_type(): def invalid_string_type(): arr = cuda.shared.array(10, dtype="int33") # noqa: F841 - with self.assertRaisesRegex(numbaError.TypingError, rgx): - cuda.jit(void())(invalid_string_type) + if HAS_NUMBA: + with self.assertRaisesRegex(NumbaTypingError, rgx): + cuda.jit(void())(invalid_string_type) + else: + with self.assertRaisesRegex(TypingError, rgx): + cuda.jit(void())(invalid_string_type) @skip_on_cudasim("Struct model array unsupported in simulator") def test_struct_model_type_static(self): diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_sm_creation.py b/numba_cuda/numba/cuda/tests/cudapy/test_sm_creation.py index 6e2f51bff..37554ec3e 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_sm_creation.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_sm_creation.py @@ -4,7 +4,12 @@ import numpy as np from numba import cuda from numba.cuda import float32, int32, void -from numba.core.errors import TypingError +from numba.cuda import HAS_NUMBA + +if HAS_NUMBA: + from numba.core.errors import TypingError +else: + from numba.cuda.core.errors import TypingError from numba.cuda.testing import unittest, CUDATestCase from numba.cuda.testing import skip_on_cudasim from .extensions_usecases import test_struct_model_type diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_ssa.py b/numba_cuda/numba/cuda/tests/cudapy/test_ssa.py index d53060d0d..2f242451f 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_ssa.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_ssa.py @@ -13,7 +13,7 @@ from numba.cuda import types from numba import cuda from numba.cuda import jit -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.extending import overload from numba.cuda.tests.support import override_config diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_typeinfer.py b/numba_cuda/numba/cuda/tests/cudapy/test_typeinfer.py index 9769c1f83..54829289b 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_typeinfer.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_typeinfer.py @@ -3,7 +3,7 @@ import itertools -from numba.core import errors +from numba.cuda.core import errors from numba.cuda import types, typing from numba.cuda.typeconv import Conversion diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_ufuncs.py b/numba_cuda/numba/cuda/tests/cudapy/test_ufuncs.py index f09205c1f..1f64b1b38 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_ufuncs.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_ufuncs.py @@ -5,9 +5,13 @@ import warnings import numpy as np import unittest +from numba.cuda import HAS_NUMBA -from numba import cuda, njit +if HAS_NUMBA: + from numba import njit +from numba import cuda from numba.cuda import config, types +from numba.cuda.testing import skip_on_standalone_numba_cuda from numba.cuda.typing.typeof import typeof from numba.cuda.np import numpy_support from numba.cuda.tests.support import TestCase @@ -65,6 +69,7 @@ def setUp(self): ), ] + @skip_on_standalone_numba_cuda @functools.lru_cache(maxsize=None) def _compile(self, pyfunc, args, nrt=False): # NOTE: to test the implementation of Numpy ufuncs, we disable diff --git a/numba_cuda/numba/cuda/tests/doc_examples/test_cpu_gpu_compat.py b/numba_cuda/numba/cuda/tests/doc_examples/test_cpu_gpu_compat.py index 6e7e838e2..1a578a9af 100644 --- a/numba_cuda/numba/cuda/tests/doc_examples/test_cpu_gpu_compat.py +++ b/numba_cuda/numba/cuda/tests/doc_examples/test_cpu_gpu_compat.py @@ -3,7 +3,11 @@ import unittest -from numba.cuda.testing import CUDATestCase, skip_on_cudasim +from numba.cuda.testing import ( + CUDATestCase, + skip_on_cudasim, + skip_on_standalone_numba_cuda, +) from numba.cuda.tests.support import captured_stdout import numpy as np @@ -25,6 +29,7 @@ def tearDown(self): self._captured_stdout.__exit__(None, None, None) super().tearDown() + @skip_on_standalone_numba_cuda def test_ex_cpu_gpu_compat(self): # ex_cpu_gpu_compat.import.begin from math import pi diff --git a/numba_cuda/numba/cuda/tests/doc_examples/test_montecarlo.py b/numba_cuda/numba/cuda/tests/doc_examples/test_montecarlo.py index 67c5d381d..67ad2bc88 100644 --- a/numba_cuda/numba/cuda/tests/doc_examples/test_montecarlo.py +++ b/numba_cuda/numba/cuda/tests/doc_examples/test_montecarlo.py @@ -3,7 +3,11 @@ import unittest -from numba.cuda.testing import CUDATestCase, skip_on_cudasim +from numba.cuda.testing import ( + CUDATestCase, + skip_on_cudasim, + skip_on_standalone_numba_cuda, +) from numba.cuda.tests.support import captured_stdout @@ -24,6 +28,7 @@ def tearDown(self): self._captured_stdout.__exit__(None, None, None) super().tearDown() + @skip_on_standalone_numba_cuda def test_ex_montecarlo(self): # ex_montecarlo.import.begin import numba diff --git a/numba_cuda/numba/cuda/tests/nocuda/test_import.py b/numba_cuda/numba/cuda/tests/nocuda/test_import.py index 7793c787e..d80ef524a 100644 --- a/numba_cuda/numba/cuda/tests/nocuda/test_import.py +++ b/numba_cuda/numba/cuda/tests/nocuda/test_import.py @@ -49,8 +49,6 @@ def test_no_impl_import(self): "numba.cuda.np.arraymath", "numba.cuda.np.npdatetime", "numba.cuda.np.npyimpl", - "numba.typed.typeddict", - "numba.typed.typedlist", ) code = "import sys; from numba import cuda; print(list(sys.modules))" diff --git a/numba_cuda/numba/cuda/tests/support.py b/numba_cuda/numba/cuda/tests/support.py index 7f8f254c2..fa5148ac2 100644 --- a/numba_cuda/numba/cuda/tests/support.py +++ b/numba_cuda/numba/cuda/tests/support.py @@ -24,7 +24,7 @@ import numpy as np from numba.cuda import types -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.core import config from numba.cuda.typing import cffi_utils from numba.cuda.memory_management.nrt import rtsys @@ -37,12 +37,13 @@ from numba.cuda.datamodel.models import OpaqueModel from numba.cuda.np import numpy_support -try: - from numba.core.extending import typeof_impl as upstream_typeof_impl +from numba.cuda import HAS_NUMBA + +if HAS_NUMBA: + from numba.core.extending import ( + typeof_impl as upstream_typeof_impl, + ) from numba.core import types as upstream_types -except ImportError: - upstream_typeof_impl = None - upstream_types = None class EnableNRTStatsMixin(object): @@ -758,7 +759,7 @@ def typeof_dummy(val, c): return dummy_type # Dual registration for cross-target tests - if upstream_typeof_impl is not None and upstream_types is not None: + if HAS_NUMBA: UpstreamDummyType = type( "DummyTypeFor{}".format(test_id), (upstream_types.Opaque,), {} ) diff --git a/numba_cuda/numba/cuda/tests/test_array_reductions.py b/numba_cuda/numba/cuda/tests/test_array_reductions.py index 06ab07a30..167095e1a 100644 --- a/numba_cuda/numba/cuda/tests/test_array_reductions.py +++ b/numba_cuda/numba/cuda/tests/test_array_reductions.py @@ -2,12 +2,14 @@ # SPDX-License-Identifier: BSD-2-Clause import numpy as np -from numba.tests.support import TestCase, MemoryLeakMixin +from numba.cuda.tests.support import TestCase, MemoryLeakMixin from numba import cuda +from numba.cuda.testing import skip_on_cudasim from numba.cuda.misc.special import literal_unroll from numba.cuda import config +@skip_on_cudasim class TestArrayReductions(MemoryLeakMixin, TestCase): """ Test array reduction methods and functions such as .sum(), .max(), etc. diff --git a/numba_cuda/numba/cuda/tests/test_extending.py b/numba_cuda/numba/cuda/tests/test_extending.py index ed24d22b7..9798de584 100644 --- a/numba_cuda/numba/cuda/tests/test_extending.py +++ b/numba_cuda/numba/cuda/tests/test_extending.py @@ -11,7 +11,13 @@ import numba from numba.cuda import jit from numba.cuda import types -from numba.core import errors +from numba.cuda import HAS_NUMBA + +if HAS_NUMBA: + from numba.core import errors +else: + from numba.cuda.core import errors + from numba.cuda.tests.support import ( TestCase, ) @@ -675,7 +681,7 @@ def create_message(func, overload_func, func_sig, ol_sig): msg.append(f" - got: {ol_sig}") lineno = inspect.getsourcelines(overload_func)[1] tmpsrcfile = inspect.getfile(overload_func) - srcfile = tmpsrcfile.replace(numba.__path__[0], "") + srcfile = tmpsrcfile.replace(numba.cuda.__path__[0], "") msg.append(f"from {srcfile}:{lineno}") msgstr = "\n" + "\n".join(msg) return msgstr diff --git a/numba_cuda/numba/cuda/tests/test_extending_types.py b/numba_cuda/numba/cuda/tests/test_extending_types.py index 08dfd4da2..60b465997 100644 --- a/numba_cuda/numba/cuda/tests/test_extending_types.py +++ b/numba_cuda/numba/cuda/tests/test_extending_types.py @@ -7,7 +7,12 @@ from numba.cuda import jit from numba.cuda import types -from numba.core.errors import TypingError, NumbaTypeError +from numba.cuda import HAS_NUMBA + +if HAS_NUMBA: + from numba.core.errors import TypingError, NumbaTypeError +else: + from numba.cuda.core.errors import TypingError, NumbaTypeError from numba.cuda.extending import make_attribute_wrapper from numba.cuda.extending import overload from numba.cuda.core.imputils import Registry diff --git a/numba_cuda/numba/cuda/types/cuda_function_type.py b/numba_cuda/numba/cuda/types/cuda_function_type.py index f4f239226..a7ef4a48b 100644 --- a/numba_cuda/numba/cuda/types/cuda_function_type.py +++ b/numba_cuda/numba/cuda/types/cuda_function_type.py @@ -13,7 +13,7 @@ from abc import ABC, abstractmethod from .abstract import Type from .. import types -from numba.core import errors +from numba.cuda.core import errors class FunctionType(Type): diff --git a/numba_cuda/numba/cuda/types/cuda_functions.py b/numba_cuda/numba/cuda/types/cuda_functions.py index da5511565..c1e3488bc 100644 --- a/numba_cuda/numba/cuda/types/cuda_functions.py +++ b/numba_cuda/numba/cuda/types/cuda_functions.py @@ -11,7 +11,7 @@ from .abstract import Callable, DTypeSpec, Dummy, Literal, Type, weakref from .common import Opaque from .misc import unliteral -from numba.core import errors +from numba.cuda.core import errors from numba.cuda import utils, types, config from numba.cuda.typeconv import Conversion diff --git a/numba_cuda/numba/cuda/typing/arraydecl.py b/numba_cuda/numba/cuda/typing/arraydecl.py index 4f322e3ec..609b437ed 100644 --- a/numba_cuda/numba/cuda/typing/arraydecl.py +++ b/numba_cuda/numba/cuda/typing/arraydecl.py @@ -518,27 +518,6 @@ def resolve_sort(self, ary, args, kws): assert not kws return signature(types.none) - @bound_function("array.argsort") - def resolve_argsort(self, ary, args, kws): - assert not args - kwargs = dict(kws) - kind = kwargs.pop("kind", types.StringLiteral("quicksort")) - if not isinstance(kind, types.StringLiteral): - raise TypingError('"kind" must be a string literal') - if kwargs: - msg = "Unsupported keywords: {!r}" - raise TypingError(msg.format([k for k in kwargs.keys()])) - if ary.ndim == 1: - - def argsort_stub(kind="quicksort"): - pass - - pysig = utils.pysignature(argsort_stub) - sig = signature(types.Array(types.intp, 1, "C"), kind).replace( - pysig=pysig - ) - return sig - @bound_function("array.view") def resolve_view(self, ary, args, kws): from .npydecl import parse_dtype diff --git a/numba_cuda/numba/cuda/typing/asnumbatype.py b/numba_cuda/numba/cuda/typing/asnumbatype.py index 50c3f9270..f6a8b9455 100644 --- a/numba_cuda/numba/cuda/typing/asnumbatype.py +++ b/numba_cuda/numba/cuda/typing/asnumbatype.py @@ -5,7 +5,7 @@ import typing as py_typing from numba.cuda.typing.typeof import typeof -from numba.core import errors +from numba.cuda.core import errors from numba.cuda import types diff --git a/numba_cuda/numba/cuda/typing/builtins.py b/numba_cuda/numba/cuda/typing/builtins.py index d8adc0781..c380cbf43 100644 --- a/numba_cuda/numba/cuda/typing/builtins.py +++ b/numba_cuda/numba/cuda/typing/builtins.py @@ -7,9 +7,7 @@ import operator from numba.cuda import types -from numba.core import errors -from numba import prange -from numba.parfors.parfor import internal_prange +from numba.cuda.core import errors from numba.cuda.typing.templates import ( AttributeTemplate, @@ -88,8 +86,6 @@ class Slice(ConcreteTemplate): @infer_global(range, typing_key=range) -@infer_global(prange, typing_key=prange) -@infer_global(internal_prange, typing_key=internal_prange) class Range(ConcreteTemplate): cases = [ signature(types.range_state32_type, types.int32), diff --git a/numba_cuda/numba/cuda/typing/collections.py b/numba_cuda/numba/cuda/typing/collections.py index a4e0dde11..8e4869d2d 100644 --- a/numba_cuda/numba/cuda/typing/collections.py +++ b/numba_cuda/numba/cuda/typing/collections.py @@ -2,7 +2,7 @@ # SPDX-License-Identifier: BSD-2-Clause from numba.cuda import types -from numba.core import errors +from numba.cuda.core import errors from numba.cuda import utils import operator from .templates import ( diff --git a/numba_cuda/numba/cuda/typing/dictdecl.py b/numba_cuda/numba/cuda/typing/dictdecl.py index 54ccc0cc0..b0a9dff1e 100644 --- a/numba_cuda/numba/cuda/typing/dictdecl.py +++ b/numba_cuda/numba/cuda/typing/dictdecl.py @@ -6,7 +6,7 @@ """ from numba.cuda import types -from numba.core import errors +from numba.cuda.core import errors from .templates import ( AbstractTemplate, Registry, diff --git a/numba_cuda/numba/cuda/typing/npdatetime.py b/numba_cuda/numba/cuda/typing/npdatetime.py index 3a72e32b5..16819f1d6 100644 --- a/numba_cuda/numba/cuda/typing/npdatetime.py +++ b/numba_cuda/numba/cuda/typing/npdatetime.py @@ -4,7 +4,7 @@ import operator from numba.cuda import types -from numba.core import errors +from numba.cuda.core import errors from numba.cuda.typing.templates import ( AbstractTemplate, Registry, diff --git a/numba_cuda/numba/cuda/typing/npydecl.py b/numba_cuda/numba/cuda/typing/npydecl.py index e5ccb069b..93eab7a44 100644 --- a/numba_cuda/numba/cuda/typing/npydecl.py +++ b/numba_cuda/numba/cuda/typing/npydecl.py @@ -504,12 +504,6 @@ def sum_stub(arr, axis, dtype): pass pysig = utils.pysignature(sum_stub) - elif self.method_name == "argsort": - - def argsort_stub(arr, kind="quicksort"): - pass - - pysig = utils.pysignature(argsort_stub) else: fmt = "numba doesn't support kwarg for {}" raise TypingError(fmt.format(self.method_name)) diff --git a/numba_cuda/numba/cuda/typing/templates.py b/numba_cuda/numba/cuda/typing/templates.py index 9c44a19f6..de4c3b464 100644 --- a/numba_cuda/numba/cuda/typing/templates.py +++ b/numba_cuda/numba/cuda/typing/templates.py @@ -25,12 +25,10 @@ from numba.cuda import utils from numba.cuda.core import targetconfig -try: - from numba.core.typing import Signature as CoreSignature +from numba.cuda import HAS_NUMBA - numba_sig_present = True -except ImportError: - numba_sig_present = False +if HAS_NUMBA: + from numba.core.typing import Signature as CoreSignature # info store for inliner callback functions e.g. cost model _inline_info = namedtuple("inline_info", "func_ir typemap calltypes signature") @@ -100,7 +98,7 @@ def __hash__(self): def __eq__(self, other): sig_types = (Signature,) - if numba_sig_present: + if HAS_NUMBA: sig_types = (Signature, CoreSignature) if isinstance(other, sig_types): return ( @@ -384,7 +382,7 @@ def apply(self, args, kws): # Enforce that *generic()* must return None or Signature if sig is not None: sig_types = (Signature,) - if numba_sig_present: + if HAS_NUMBA: sig_types = (Signature, CoreSignature) if not isinstance(sig, sig_types): raise AssertionError( @@ -409,7 +407,9 @@ def unpack_opt(x): def get_template_info(self): impl = getattr(self, "generic") - basepath = os.path.dirname(os.path.dirname(numba.__file__)) + basepath = os.path.dirname( + os.path.dirname(os.path.dirname(numba.cuda.__file__)) + ) code, firstlineno, path = self.get_source_code_info(impl) sig = str(utils.pysignature(impl)) @@ -496,7 +496,9 @@ def unpack_opt(x): def get_template_info(self): impl = getattr(self, "generic") - basepath = os.path.dirname(os.path.dirname(numba.__file__)) + basepath = os.path.dirname( + os.path.dirname(os.path.dirname(numba.cuda.__file__)) + ) code, firstlineno, path = self.get_source_code_info(impl) sig = str(utils.pysignature(impl)) info = { @@ -885,7 +887,9 @@ def get_source_info(cls): - "docstring": str The docstring of the definition. """ - basepath = os.path.dirname(os.path.dirname(numba.__file__)) + basepath = os.path.dirname( + os.path.dirname(os.path.dirname(numba.cuda.__file__)) + ) impl = cls._overload_func code, firstlineno, path = cls.get_source_code_info(impl) sig = str(utils.pysignature(impl)) @@ -900,7 +904,9 @@ def get_source_info(cls): return info def get_template_info(self): - basepath = os.path.dirname(os.path.dirname(numba.__file__)) + basepath = os.path.dirname( + os.path.dirname(os.path.dirname(numba.cuda.__file__)) + ) impl = self._overload_func code, firstlineno, path = self.get_source_code_info(impl) sig = str(utils.pysignature(impl)) @@ -1057,7 +1063,9 @@ def get_impl_key(self, sig): return self._overload_cache[sig.args] def get_template_info(self): - basepath = os.path.dirname(os.path.dirname(numba.__file__)) + basepath = os.path.dirname( + os.path.dirname(os.path.dirname(numba.cuda.__file__)) + ) impl = self._definition_func code, firstlineno, path = self.get_source_code_info(impl) sig = str(utils.pysignature(impl)) @@ -1219,7 +1227,9 @@ def generic(_, args, kws): return sig.as_method() def get_template_info(self): - basepath = os.path.dirname(os.path.dirname(numba.__file__)) + basepath = os.path.dirname( + os.path.dirname(os.path.dirname(numba.cuda.__file__)) + ) impl = self._overload_func code, firstlineno, path = self.get_source_code_info(impl) sig = str(utils.pysignature(impl)) diff --git a/numba_cuda/numba/cuda/typing/typeof.py b/numba_cuda/numba/cuda/typing/typeof.py index f199c5af4..a3091d282 100644 --- a/numba_cuda/numba/cuda/typing/typeof.py +++ b/numba_cuda/numba/cuda/typing/typeof.py @@ -8,9 +8,9 @@ import numpy as np from numpy.random.bit_generator import BitGenerator - +from numba.cuda import HAS_NUMBA from numba.cuda import types -from numba.core import errors +from numba.cuda.core import errors from numba.cuda import utils from numba.cuda.np import numpy_support @@ -65,15 +65,13 @@ def typeof_impl(val, c): if cffi_utils.is_ffi_instance(val): return types.ffi - try: + if HAS_NUMBA: # Fallback to Numba's typeof_impl for third-party registrations from numba.core.typing.typeof import typeof_impl as core_typeof_impl tp = core_typeof_impl(val, c) if tp is not None: return tp - except Exception: - pass return None @@ -119,16 +117,6 @@ def _typeof_type(val, c): if issubclass(val, types.Type): return types.TypeRef(val) - from numba.typed import Dict - - if issubclass(val, Dict): - return types.TypeRef(types.DictType) - - from numba.typed import List - - if issubclass(val, List): - return types.TypeRef(types.ListType) - @typeof_impl.register(bool) def _typeof_bool(val, c):