Merge remote-tracking branch 'NVIDIA/main' into fix-target-registration

Author: gmarkall

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/_internal/cuda_fp16.py

@@ -21,7 +21,7 @@
 from numba.cuda import types
 from numba.cuda.cudadrv.driver import _have_nvjitlink
 from numba.cuda.datamodel import PrimitiveModel, StructModel
-from numba.core.errors import NumbaPerformanceWarning
+from numba.cuda.core.errors import NumbaPerformanceWarning
 from numba.cuda.extending import (
     lower_cast,
     make_attribute_wrapper,

numba_cuda/numba/cuda/compiler.py

@@ -6,12 +6,12 @@
 from warnings import warn, catch_warnings, simplefilter
 import copy
 
-from numba.core import ir as numba_ir
-from numba.core import bytecode
+from numba.cuda.core import ir as numba_ir
+from numba.cuda.core import bytecode
 from numba.cuda import types
 from numba.cuda.core.options import ParallelOptions
 from numba.core.compiler_lock import global_compiler_lock
-from numba.core.errors import NumbaWarning, NumbaInvalidConfigWarning
+from numba.cuda.core.errors import NumbaWarning, NumbaInvalidConfigWarning
 from numba.cuda.core.interpreter import Interpreter
 
 from numba.cuda import cgutils, typing, lowering, nvvmutils, utils
@@ -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

@@ -3,13 +3,14 @@
 
 from collections import namedtuple, defaultdict
 from numba.cuda import types
-from numba.core import ir, errors
+from numba.cuda.core import ir
+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
@@ -353,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:
@@ -382,7 +383,7 @@ def find_literally_calls(func_ir, argtypes):
     argtypes : Sequence[numba.types.Type]
         The argument types.
     """
-    from numba.core import ir_utils
+    from numba.cuda.core import ir_utils
 
     marked_args = set()
     first_loc = {}

numba_cuda/numba/cuda/core/annotations/type_annotations.py

@@ -11,7 +11,7 @@
 import textwrap
 from io import StringIO
 
-from numba.core import ir
+from numba.cuda.core import ir
 
 
 class SourceLines(Mapping):

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/boxing.py

@@ -9,7 +9,7 @@
 
 from numba.cuda import types, cgutils
 from numba.cuda.core.pythonapi import box, unbox, reflect, NativeValue
-from numba.core.errors import NumbaNotImplementedError, TypingError
+from numba.cuda.core.errors import NumbaNotImplementedError, TypingError
 from numba.cuda.typing.typeof import typeof, Purpose
 
 from numba.cuda.cpython import setobj, listobj

numba_cuda/numba/cuda/core/bytecode.py

@@ -8,7 +8,8 @@
 
 from types import CodeType, ModuleType
 
-from numba.core import errors, 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,8 +3,8 @@
 
 from numba.cuda.core.tracing import event
 
-from numba.core import errors
-from numba.core.errors import CompilerError
+from numba.cuda.core import errors
+from numba.cuda.core.errors import CompilerError
 
 from numba.cuda.core import callconv, config, bytecode
 from numba.cuda.core.untyped_passes import ExtractByteCode, FixupArgs