[Refactor][NFC][Cleanups] Update imports to upstream numba to use the numba.cuda modules

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 (

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"
 

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:

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 (

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
 

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
 
 

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

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

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)
 

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
 

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:

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)

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
 
 

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)

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 = (

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 (