[Refactor][NFC] Vendor in numba.misc for CUDA-specific changes

numba_cuda/numba/cuda/core/caching.py

@@ -14,7 +14,7 @@
 import tempfile
 import sys
 
-from numba.misc.appdirs import AppDirs
+from numba.cuda.misc.appdirs import AppDirs
 from pathlib import Path
 
 from numba.cuda.core import config

numba_cuda/numba/cuda/core/codegen.py

@@ -2,7 +2,7 @@
 # SPDX-License-Identifier: BSD-2-Clause
 
 from abc import abstractmethod, ABCMeta
-from numba.misc.llvm_pass_timings import PassTimingsCollection
+from numba.cuda.misc.llvm_pass_timings import PassTimingsCollection
 
 
 class CodeLibrary(metaclass=ABCMeta):

numba_cuda/numba/cuda/core/inline_closurecall.py

@@ -41,7 +41,7 @@
 from numba.np.unsafe.ndarray import empty_inferred as unsafe_empty_inferred
 import numpy as np
 import operator
-import numba.misc.special
+from numba.cuda.misc.special import prange
 
 """
 Variable enable_inline_arraycall is only used for testing purpose.
@@ -1054,10 +1054,7 @@ def _find_iter_range(func_ir, range_iter_var, swapped):
     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 == numba.misc.special.prange
-        )
+        and (func_def.value is range or func_def.value == prange)
     )
     nargs = len(range_def.args)
     swapping = [('"array comprehension"', "closure of"), range_def.func.loc]

numba_cuda/numba/cuda/core/typeinfer.py

@@ -1692,7 +1692,7 @@ def typeof_global(self, inst, target, gvar):
                 # as a global variable
                 typ = types.Dispatcher(_temporary_dispatcher_map[gvar.name])
             else:
-                from numba.misc import special
+                from numba.cuda.misc import special
 
                 nm = gvar.name
                 # check if the problem is actually a name error

numba_cuda/numba/cuda/core/untyped_passes.py

@@ -22,7 +22,7 @@
 from numba.cuda.core.interpreter import Interpreter
 
 
-from numba.misc.special import literal_unroll
+from numba.cuda.misc.special import literal_unroll
 from numba.cuda.core.analysis import dead_branch_prune
 from numba.core.analysis import (
     rewrite_semantic_constants,

numba_cuda/numba/cuda/lowering.py

@@ -29,13 +29,13 @@
 from numba.cuda.core.funcdesc import default_mangler
 from numba.cuda.core.environment import Environment
 from numba.core.analysis import compute_use_defs, must_use_alloca
-from numba.misc.firstlinefinder import get_func_body_first_lineno
+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.misc.coverage_support import get_registered_loc_notify
+    from numba.cuda.misc.coverage_support import get_registered_loc_notify
 
 
 _VarArgItem = namedtuple("_VarArgItem", ("vararg", "index"))