[Refactor][NFC][Cleanups] Update imports to upstream numba to use the numba.cuda modules (#561)

This PR updates the vast majority of imports to upstream numba modules,
that had already been vendored in for future CUDA-specific changes.
After this PR, there should be exactly 3 (not strictly necessary)
upstream numba imports remaining - one in testing vectorization for in
test_math, one in random.py, and the global compiler_lock.

It seems from the existing comments that removing the CPU jit in
random.py would result in considerable performance drops, so I've just
added a guard for numba being available.

Some lowering / overload implementations that depend on
numba.core.misc.{mergesort,quicksort} have been removed. These
are presently not regarded as supported, and are unlikely to be
performant implementations in a CUDA-specific context.

There are many tests that rely on CPU jitting, these tests are now
guarded to not run in environments where the numba package is not
available. This is done using a global importlib check in numba-cuda's
`__init__.py`. This cannot be moved into numba.cuda.config since the
upstream numba doesn't have this flag, and we intend on forwarding calls
to upstream numba when it is available. All previously sporadic guards
for this *should* have been updated to use this global flag now.

Support for prange, pndindex, and stencils has been removed, likewise
for typed.Dict, and typed.List.

 Also, this removes all mention of the legacy type
system flag since we implicitly assume the legacy type system in place.

Author: atmnp

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