From 94014ce052c0798c078332e696c522a5dfd803a5 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Tue, 28 Oct 2025 14:47:02 -0700
Subject: [PATCH 01/25] [Refactor][NFC] Remove references to numba.__file__ and
 numba.__path__

---
 numba_cuda/numba/cuda/tests/test_extending.py |  2 +-
 numba_cuda/numba/cuda/typing/templates.py     | 12 ++++++------
 2 files changed, 7 insertions(+), 7 deletions(-)

diff --git a/numba_cuda/numba/cuda/tests/test_extending.py b/numba_cuda/numba/cuda/tests/test_extending.py
index ed24d22b7..d26a07957 100644
--- a/numba_cuda/numba/cuda/tests/test_extending.py
+++ b/numba_cuda/numba/cuda/tests/test_extending.py
@@ -675,7 +675,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/typing/templates.py b/numba_cuda/numba/cuda/typing/templates.py
index 9c44a19f6..235bc758f 100644
--- a/numba_cuda/numba/cuda/typing/templates.py
+++ b/numba_cuda/numba/cuda/typing/templates.py
@@ -409,7 +409,7 @@ 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(numba.cuda.__file__))
 
         code, firstlineno, path = self.get_source_code_info(impl)
         sig = str(utils.pysignature(impl))
@@ -496,7 +496,7 @@ 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(numba.cuda.__file__))
         code, firstlineno, path = self.get_source_code_info(impl)
         sig = str(utils.pysignature(impl))
         info = {
@@ -885,7 +885,7 @@ 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(numba.cuda.__file__))
         impl = cls._overload_func
         code, firstlineno, path = cls.get_source_code_info(impl)
         sig = str(utils.pysignature(impl))
@@ -900,7 +900,7 @@ 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(numba.cuda.__file__))
         impl = self._overload_func
         code, firstlineno, path = self.get_source_code_info(impl)
         sig = str(utils.pysignature(impl))
@@ -1057,7 +1057,7 @@ 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(numba.cuda.__file__))
         impl = self._definition_func
         code, firstlineno, path = self.get_source_code_info(impl)
         sig = str(utils.pysignature(impl))
@@ -1219,7 +1219,7 @@ 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(numba.cuda.__file__))
                 impl = self._overload_func
                 code, firstlineno, path = self.get_source_code_info(impl)
                 sig = str(utils.pysignature(impl))

From ab627e828f27aeee8ae9b886f7a4d65bf64d63e3 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Tue, 28 Oct 2025 14:52:52 -0700
Subject: [PATCH 02/25] [Refactor][NFC] Remove references to pndindex,
 literal_unroll, stencils

---
 numba_cuda/numba/cuda/__init__.py             |  1 +
 .../numba/cuda/core/inline_closurecall.py     | 91 -------------------
 numba_cuda/numba/cuda/core/ir_utils.py        | 10 +-
 numba_cuda/numba/cuda/np/arrayobj.py          |  4 +-
 .../np/polynomial/polynomial_functions.py     |  2 +-
 5 files changed, 7 insertions(+), 101 deletions(-)

diff --git a/numba_cuda/numba/cuda/__init__.py b/numba_cuda/numba/cuda/__init__.py
index ad92b20ac..f259dfef3 100644
--- a/numba_cuda/numba/cuda/__init__.py
+++ b/numba_cuda/numba/cuda/__init__.py
@@ -66,3 +66,4 @@
     )
 
 from numba.cuda.np.ufunc import vectorize, guvectorize
+from numba.cuda.misc.special import literal_unroll
diff --git a/numba_cuda/numba/cuda/core/inline_closurecall.py b/numba_cuda/numba/cuda/core/inline_closurecall.py
index c45f3265a..5aa6c0660 100644
--- a/numba_cuda/numba/cuda/core/inline_closurecall.py
+++ b/numba_cuda/numba/cuda/core/inline_closurecall.py
@@ -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"
diff --git a/numba_cuda/numba/cuda/core/ir_utils.py b/numba_cuda/numba/cuda/core/ir_utils.py
index c290cda12..ebe9f348a 100644
--- a/numba_cuda/numba/cuda/core/ir_utils.py
+++ b/numba_cuda/numba/cuda/core/ir_utils.py
@@ -9,7 +9,7 @@
 import collections
 import warnings
 
-import numba
+import numba.cuda
 from numba.cuda import types
 from numba.cuda.core import ir
 from numba.cuda import typing
@@ -816,10 +816,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]
@@ -1992,7 +1990,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 +2331,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 +2489,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/np/arrayobj.py b/numba_cuda/numba/cuda/np/arrayobj.py
index fcf979f4f..e8f7c9c4e 100644
--- a/numba_cuda/numba/cuda/np/arrayobj.py
+++ b/numba_cuda/numba/cuda/np/arrayobj.py
@@ -16,7 +16,7 @@
 
 import numpy as np
 
-from numba import pndindex, literal_unroll
+from numba.cuda import literal_unroll
 from numba.cuda import types, typing
 from numba.core import errors
 from numba.cuda import cgutils, extending
@@ -4529,7 +4529,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 +4545,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)"""
diff --git a/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py b/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py
index 4cd48baf0..6b57f4ef8 100644
--- a/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py
+++ b/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py
@@ -9,7 +9,7 @@
 from numpy.polynomial import polynomial as poly
 from numpy.polynomial import polyutils as pu
 
-from numba import literal_unroll
+from numba.cuda import literal_unroll
 from numba.cuda import types
 from numba.core import errors
 from numba.cuda.extending import overload

From 679d26f428b9818b3137af061d9e9f15ab8767a6 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Tue, 28 Oct 2025 14:58:10 -0700
Subject: [PATCH 03/25] [Refactor][NFC] Remove reference to numba.__version__

---
 numba_cuda/numba/cuda/__init__.py    | 3 +++
 numba_cuda/numba/cuda/core/errors.py | 4 +++-
 2 files changed, 6 insertions(+), 1 deletion(-)

diff --git a/numba_cuda/numba/cuda/__init__.py b/numba_cuda/numba/cuda/__init__.py
index f259dfef3..97bcfdc3f 100644
--- a/numba_cuda/numba/cuda/__init__.py
+++ b/numba_cuda/numba/cuda/__init__.py
@@ -7,6 +7,9 @@
 import warnings
 import sys
 
+# Import version from the parent package
+from numba_cuda._version import __version__
+
 # Re-export types itself
 import numba.cuda.types as types
 
diff --git a/numba_cuda/numba/cuda/core/errors.py b/numba_cuda/numba/cuda/core/errors.py
index aef7b26e5..e2303263d 100644
--- a/numba_cuda/numba/cuda/core/errors.py
+++ b/numba_cuda/numba/cuda/core/errors.py
@@ -15,6 +15,8 @@
 from functools import wraps
 from abc import abstractmethod
 
+from numba_cuda._version import __version__
+
 # Filled at the end
 __all__ = []
 
@@ -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)
+""" % (__version__, feedback_details)
 
 error_extras = dict()
 error_extras["unsupported_error"] = unsupported_error_info

From f99f2069438b6243c0fde41e2307c3212560eafc Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Tue, 28 Oct 2025 15:03:22 -0700
Subject: [PATCH 04/25] [Refactor][NFC] Vendor-in mergesort and quicksort for
 future CUDA-specific changes

---
 numba_cuda/numba/cuda/__init__.py        |   1 +
 numba_cuda/numba/cuda/cpython/listobj.py |   2 +-
 numba_cuda/numba/cuda/misc/mergesort.py  | 136 +++++++++++
 numba_cuda/numba/cuda/misc/quicksort.py  | 273 +++++++++++++++++++++++
 numba_cuda/numba/cuda/np/arrayobj.py     |   2 +-
 5 files changed, 412 insertions(+), 2 deletions(-)
 create mode 100644 numba_cuda/numba/cuda/misc/mergesort.py
 create mode 100644 numba_cuda/numba/cuda/misc/quicksort.py

diff --git a/numba_cuda/numba/cuda/__init__.py b/numba_cuda/numba/cuda/__init__.py
index 97bcfdc3f..04793843c 100644
--- a/numba_cuda/numba/cuda/__init__.py
+++ b/numba_cuda/numba/cuda/__init__.py
@@ -69,4 +69,5 @@
     )
 
 from numba.cuda.np.ufunc import vectorize, guvectorize
+from numba.cuda.misc import quicksort, mergesort
 from numba.cuda.misc.special import literal_unroll
diff --git a/numba_cuda/numba/cuda/cpython/listobj.py b/numba_cuda/numba/cuda/cpython/listobj.py
index de93b23ee..b1904e514 100644
--- a/numba_cuda/numba/cuda/cpython/listobj.py
+++ b/numba_cuda/numba/cuda/cpython/listobj.py
@@ -19,7 +19,7 @@
     RefType,
 )
 from numba.cuda.extending import overload_method, overload
-from numba.misc import quicksort
+from numba.cuda.misc import quicksort
 from numba.cuda.cpython import slicing
 from numba import literal_unroll
 
diff --git a/numba_cuda/numba/cuda/misc/mergesort.py b/numba_cuda/numba/cuda/misc/mergesort.py
new file mode 100644
index 000000000..79758485f
--- /dev/null
+++ b/numba_cuda/numba/cuda/misc/mergesort.py
@@ -0,0 +1,136 @@
+# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: BSD-2-Clause
+
+"""
+The same algorithm as translated from numpy.
+See numpy/core/src/npysort/mergesort.c.src.
+The high-level numba code is adding a little overhead comparing to
+the pure-C implementation in numpy.
+"""
+
+import numpy as np
+from collections import namedtuple
+
+# Array size smaller than this will be sorted by insertion sort
+SMALL_MERGESORT = 20
+
+
+MergesortImplementation = namedtuple(
+    "MergesortImplementation",
+    [
+        "run_mergesort",
+    ],
+)
+
+
+def make_mergesort_impl(wrap, lt=None, is_argsort=False):
+    kwargs_lite = dict(no_cpython_wrapper=True, _nrt=False)
+
+    # The less than
+    if lt is None:
+
+        @wrap(**kwargs_lite)
+        def lt(a, b):
+            return a < b
+    else:
+        lt = wrap(**kwargs_lite)(lt)
+
+    if is_argsort:
+
+        @wrap(**kwargs_lite)
+        def lessthan(a, b, vals):
+            return lt(vals[a], vals[b])
+    else:
+
+        @wrap(**kwargs_lite)
+        def lessthan(a, b, vals):
+            return lt(a, b)
+
+    @wrap(**kwargs_lite)
+    def argmergesort_inner(arr, vals, ws):
+        """The actual mergesort function
+
+        Parameters
+        ----------
+        arr : array [read+write]
+            The values being sorted inplace.  For argsort, this is the
+            indices.
+        vals : array [readonly]
+            ``None`` for normal sort.  In argsort, this is the actual array values.
+        ws : array [write]
+            The workspace.  Must be of size ``arr.size // 2``
+        """
+        if arr.size > SMALL_MERGESORT:
+            # Merge sort
+            mid = arr.size // 2
+
+            argmergesort_inner(arr[:mid], vals, ws)
+            argmergesort_inner(arr[mid:], vals, ws)
+
+            # Copy left half into workspace so we don't overwrite it
+            for i in range(mid):
+                ws[i] = arr[i]
+
+            # Merge
+            left = ws[:mid]
+            right = arr[mid:]
+            out = arr
+
+            i = j = k = 0
+            while i < left.size and j < right.size:
+                if not lessthan(right[j], left[i], vals):
+                    out[k] = left[i]
+                    i += 1
+                else:
+                    out[k] = right[j]
+                    j += 1
+                k += 1
+
+            # Leftovers
+            while i < left.size:
+                out[k] = left[i]
+                i += 1
+                k += 1
+
+            while j < right.size:
+                out[k] = right[j]
+                j += 1
+                k += 1
+        else:
+            # Insertion sort
+            i = 1
+            while i < arr.size:
+                j = i
+                while j > 0 and lessthan(arr[j], arr[j - 1], vals):
+                    arr[j - 1], arr[j] = arr[j], arr[j - 1]
+                    j -= 1
+                i += 1
+
+    # The top-level entry points
+
+    @wrap(no_cpython_wrapper=True)
+    def mergesort(arr):
+        "Inplace"
+        ws = np.empty(arr.size // 2, dtype=arr.dtype)
+        argmergesort_inner(arr, None, ws)
+        return arr
+
+    @wrap(no_cpython_wrapper=True)
+    def argmergesort(arr):
+        "Out-of-place"
+        idxs = np.arange(arr.size)
+        ws = np.empty(arr.size // 2, dtype=idxs.dtype)
+        argmergesort_inner(idxs, arr, ws)
+        return idxs
+
+    return MergesortImplementation(
+        run_mergesort=(argmergesort if is_argsort else mergesort)
+    )
+
+
+def make_jit_mergesort(*args, **kwargs):
+    from numba import njit
+
+    # NOTE: wrap with njit to allow recursion
+    #       because @register_jitable => @overload doesn't support recursion
+    return make_mergesort_impl(njit, *args, **kwargs)
diff --git a/numba_cuda/numba/cuda/misc/quicksort.py b/numba_cuda/numba/cuda/misc/quicksort.py
new file mode 100644
index 000000000..042fc2aba
--- /dev/null
+++ b/numba_cuda/numba/cuda/misc/quicksort.py
@@ -0,0 +1,273 @@
+# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: BSD-2-Clause
+
+import collections
+
+import numpy as np
+
+from numba.cuda import types, config
+
+
+QuicksortImplementation = collections.namedtuple(
+    "QuicksortImplementation",
+    (  # The compile function itself
+        "compile",
+        # All subroutines exercised by test_sort
+        "partition",
+        "partition3",
+        "insertion_sort",
+        # The top-level function
+        "run_quicksort",
+    ),
+)
+
+
+Partition = collections.namedtuple("Partition", ("start", "stop"))
+
+# Under this size, switch to a simple insertion sort
+SMALL_QUICKSORT = 15
+
+MAX_STACK = 100
+
+
+def make_quicksort_impl(
+    wrap, lt=None, is_argsort=False, is_list=False, is_np_array=False
+):
+    if config.USE_LEGACY_TYPE_SYSTEM:
+        intp = types.intp
+    else:
+        intp = types.py_int
+    zero = intp(0)
+
+    # Two subroutines to make the core algorithm generic wrt. argsort
+    # or normal sorting.  Note the genericity may make basic sort()
+    # slightly slower (~5%)
+    if is_argsort:
+        if is_list:
+
+            @wrap
+            def make_res(A):
+                return [x for x in range(len(A))]
+        else:
+
+            @wrap
+            def make_res(A):
+                return np.arange(A.size)
+
+        @wrap
+        def GET(A, idx_or_val):
+            return A[idx_or_val]
+
+    else:
+
+        @wrap
+        def make_res(A):
+            return A
+
+        @wrap
+        def GET(A, idx_or_val):
+            return idx_or_val
+
+    def default_lt(a, b):
+        """
+        Trivial comparison function between two keys.
+        """
+        return a < b
+
+    LT = wrap(lt if lt is not None else default_lt)
+
+    @wrap
+    def insertion_sort(A, R, low, high):
+        """
+        Insertion sort A[low:high + 1]. Note the inclusive bounds.
+        """
+        assert low >= 0
+        if high <= low:
+            return
+
+        for i in range(low + 1, high + 1):
+            k = R[i]
+            v = GET(A, k)
+            # Insert v into A[low:i]
+            j = i
+            while j > low and LT(v, GET(A, R[j - 1])):
+                # Make place for moving A[i] downwards
+                R[j] = R[j - 1]
+                j -= 1
+            R[j] = k
+
+    @wrap
+    def partition(A, R, low, high):
+        """
+        Partition A[low:high + 1] around a chosen pivot.  The pivot's index
+        is returned.
+        """
+        assert low >= 0
+        assert high > low
+
+        mid = (low + high) >> 1
+        # NOTE: the pattern of swaps below for the pivot choice and the
+        # partitioning gives good results (i.e. regular O(n log n))
+        # on sorted, reverse-sorted, and uniform arrays.  Subtle changes
+        # risk breaking this property.
+
+        # median of three {low, middle, high}
+        if LT(GET(A, R[mid]), GET(A, R[low])):
+            R[low], R[mid] = R[mid], R[low]
+        if LT(GET(A, R[high]), GET(A, R[mid])):
+            R[high], R[mid] = R[mid], R[high]
+        if LT(GET(A, R[mid]), GET(A, R[low])):
+            R[low], R[mid] = R[mid], R[low]
+        pivot = GET(A, R[mid])
+
+        # Temporarily stash the pivot at the end
+        R[high], R[mid] = R[mid], R[high]
+        i = low
+        j = high - 1
+        while True:
+            while i < high and LT(GET(A, R[i]), pivot):
+                i += 1
+            while j >= low and LT(pivot, GET(A, R[j])):
+                j -= 1
+            if i >= j:
+                break
+            R[i], R[j] = R[j], R[i]
+            i += 1
+            j -= 1
+        # Put the pivot back in its final place (all items before `i`
+        # are smaller than the pivot, all items at/after `i` are larger)
+        R[i], R[high] = R[high], R[i]
+        return i
+
+    @wrap
+    def partition3(A, low, high):
+        """
+        Three-way partition [low, high) around a chosen pivot.
+        A tuple (lt, gt) is returned such that:
+            - all elements in [low, lt) are < pivot
+            - all elements in [lt, gt] are == pivot
+            - all elements in (gt, high] are > pivot
+        """
+        mid = (low + high) >> 1
+        # median of three {low, middle, high}
+        if LT(A[mid], A[low]):
+            A[low], A[mid] = A[mid], A[low]
+        if LT(A[high], A[mid]):
+            A[high], A[mid] = A[mid], A[high]
+        if LT(A[mid], A[low]):
+            A[low], A[mid] = A[mid], A[low]
+        pivot = A[mid]
+
+        A[low], A[mid] = A[mid], A[low]
+        lt = low
+        gt = high
+        i = low + 1
+        while i <= gt:
+            if LT(A[i], pivot):
+                A[lt], A[i] = A[i], A[lt]
+                lt += 1
+                i += 1
+            elif LT(pivot, A[i]):
+                A[gt], A[i] = A[i], A[gt]
+                gt -= 1
+            else:
+                i += 1
+        return lt, gt
+
+    @wrap
+    def run_quicksort1(A):
+        R = make_res(A)
+
+        if len(A) < 2:
+            return R
+
+        stack = [Partition(zero, zero)] * MAX_STACK
+        stack[0] = Partition(zero, len(A) - 1)
+        n = 1
+
+        while n > 0:
+            n -= 1
+            low, high = stack[n]
+            # Partition until it becomes more efficient to do an insertion sort
+            while high - low >= SMALL_QUICKSORT:
+                assert n < MAX_STACK
+                i = partition(A, R, low, high)
+                # Push largest partition on the stack
+                if high - i > i - low:
+                    # Right is larger
+                    if high > i:
+                        stack[n] = Partition(i + 1, high)
+                        n += 1
+                    high = i - 1
+                else:
+                    if i > low:
+                        stack[n] = Partition(low, i - 1)
+                        n += 1
+                    low = i + 1
+
+            insertion_sort(A, R, low, high)
+
+        return R
+
+    if is_np_array:
+
+        @wrap
+        def run_quicksort(A):
+            if A.ndim == 1:
+                return run_quicksort1(A)
+            else:
+                for idx in np.ndindex(A.shape[:-1]):
+                    run_quicksort1(A[idx])
+                return A
+    else:
+
+        @wrap
+        def run_quicksort(A):
+            return run_quicksort1(A)
+
+    # Unused quicksort implementation based on 3-way partitioning; the
+    # partitioning scheme turns out exhibiting bad behaviour on sorted arrays.
+    @wrap
+    def _run_quicksort(A):
+        stack = [Partition(zero, zero)] * 100
+        stack[0] = Partition(zero, len(A) - 1)
+        n = 1
+
+        while n > 0:
+            n -= 1
+            low, high = stack[n]
+            # Partition until it becomes more efficient to do an insertion sort
+            while high - low >= SMALL_QUICKSORT:
+                assert n < MAX_STACK
+                l, r = partition3(A, low, high)
+                # One trivial (empty) partition => iterate on the other
+                if r == high:
+                    high = l - 1
+                elif l == low:
+                    low = r + 1
+                # Push largest partition on the stack
+                elif high - r > l - low:
+                    # Right is larger
+                    stack[n] = Partition(r + 1, high)
+                    n += 1
+                    high = l - 1
+                else:
+                    stack[n] = Partition(low, l - 1)
+                    n += 1
+                    low = r + 1
+
+            insertion_sort(A, low, high)
+
+    return QuicksortImplementation(
+        wrap, partition, partition3, insertion_sort, run_quicksort
+    )
+
+
+def make_py_quicksort(*args, **kwargs):
+    return make_quicksort_impl((lambda f: f), *args, **kwargs)
+
+
+def make_jit_quicksort(*args, **kwargs):
+    from numba.core.extending import register_jitable
+
+    return make_quicksort_impl((lambda f: register_jitable(f)), *args, **kwargs)
diff --git a/numba_cuda/numba/cuda/np/arrayobj.py b/numba_cuda/numba/cuda/np/arrayobj.py
index e8f7c9c4e..6df416667 100644
--- a/numba_cuda/numba/cuda/np/arrayobj.py
+++ b/numba_cuda/numba/cuda/np/arrayobj.py
@@ -54,7 +54,7 @@
     intrinsic,
     overload_attribute,
 )
-from numba.misc import quicksort, mergesort
+from numba.cuda.misc import quicksort, mergesort
 from numba.cuda.cpython import slicing
 from numba.cuda.cpython.unsafe.tuple import (
     tuple_setitem,

From c6ca4363c7150404c0e2fdfdbdf648aeaff34042 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Tue, 28 Oct 2025 15:05:07 -0700
Subject: [PATCH 05/25] [Refactor][NFC] Remove references to prange

---
 numba_cuda/numba/cuda/typing/builtins.py | 2 --
 1 file changed, 2 deletions(-)

diff --git a/numba_cuda/numba/cuda/typing/builtins.py b/numba_cuda/numba/cuda/typing/builtins.py
index d8adc0781..82b1deeaf 100644
--- a/numba_cuda/numba/cuda/typing/builtins.py
+++ b/numba_cuda/numba/cuda/typing/builtins.py
@@ -9,7 +9,6 @@
 from numba.cuda import types
 from numba.core import errors
 from numba import prange
-from numba.parfors.parfor import internal_prange
 
 from numba.cuda.typing.templates import (
     AttributeTemplate,
@@ -89,7 +88,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),

From b1f91abe0d525955f1bed21b867c022406ef23f2 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Tue, 28 Oct 2025 15:06:52 -0700
Subject: [PATCH 06/25] [Refactor][NFC] Remove references to numba.misc, update
 to numba.cuda.misc

---
 numba_cuda/numba/cuda/core/analysis.py    | 2 +-
 numba_cuda/numba/cuda/core/ir.py          | 4 ++--
 numba_cuda/numba/cuda/cpython/builtins.py | 2 +-
 3 files changed, 4 insertions(+), 4 deletions(-)

diff --git a/numba_cuda/numba/cuda/core/analysis.py b/numba_cuda/numba/cuda/core/analysis.py
index 4cc19cc37..d998cc34b 100644
--- a/numba_cuda/numba/cuda/core/analysis.py
+++ b/numba_cuda/numba/cuda/core/analysis.py
@@ -10,7 +10,7 @@
 from functools import reduce
 
 from .controlflow import CFGraph
-from numba.misc import special
+from numba.cuda.misc import special
 
 #
 # Analysis related to variable lifetime
diff --git a/numba_cuda/numba/cuda/core/ir.py b/numba_cuda/numba/cuda/core/ir.py
index 9951fdfbe..281085e32 100644
--- a/numba_cuda/numba/cuda/core/ir.py
+++ b/numba_cuda/numba/cuda/core/ir.py
@@ -1708,8 +1708,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/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

From eb1accdf0fda7a3900481e747556ac8d86cc0e28 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Tue, 28 Oct 2025 15:08:43 -0700
Subject: [PATCH 07/25] [Refactor][NFC] Remove more prange references

---
 numba_cuda/numba/cuda/core/inline_closurecall.py | 6 +-----
 numba_cuda/numba/cuda/typing/builtins.py         | 2 --
 2 files changed, 1 insertion(+), 7 deletions(-)

diff --git a/numba_cuda/numba/cuda/core/inline_closurecall.py b/numba_cuda/numba/cuda/core/inline_closurecall.py
index 5aa6c0660..db9f0eb08 100644
--- a/numba_cuda/numba/cuda/core/inline_closurecall.py
+++ b/numba_cuda/numba/cuda/core/inline_closurecall.py
@@ -43,7 +43,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.
@@ -969,10 +968,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:
diff --git a/numba_cuda/numba/cuda/typing/builtins.py b/numba_cuda/numba/cuda/typing/builtins.py
index 82b1deeaf..70b7ef532 100644
--- a/numba_cuda/numba/cuda/typing/builtins.py
+++ b/numba_cuda/numba/cuda/typing/builtins.py
@@ -8,7 +8,6 @@
 
 from numba.cuda import types
 from numba.core import errors
-from numba import prange
 
 from numba.cuda.typing.templates import (
     AttributeTemplate,
@@ -87,7 +86,6 @@ class Slice(ConcreteTemplate):
 
 
 @infer_global(range, typing_key=range)
-@infer_global(prange, typing_key=prange)
 class Range(ConcreteTemplate):
     cases = [
         signature(types.range_state32_type, types.int32),

From 0782fba60b33e0c183c03a97ec90f82c3c88804b Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Tue, 28 Oct 2025 15:30:38 -0700
Subject: [PATCH 08/25] [Refactor][NFC] Update references to numba.errors to
 now reference numba.cuda.errors

---
 numba_cuda/numba/cuda/__init__.py                         | 7 +++++++
 numba_cuda/numba/cuda/core/analysis.py                    | 2 +-
 numba_cuda/numba/cuda/core/base.py                        | 2 +-
 numba_cuda/numba/cuda/core/bytecode.py                    | 2 +-
 numba_cuda/numba/cuda/core/compiler.py                    | 2 +-
 numba_cuda/numba/cuda/core/compiler_machinery.py          | 2 +-
 numba_cuda/numba/cuda/core/inline_closurecall.py          | 2 +-
 numba_cuda/numba/cuda/core/interpreter.py                 | 2 +-
 numba_cuda/numba/cuda/core/ir.py                          | 2 +-
 numba_cuda/numba/cuda/core/pythonapi.py                   | 2 +-
 numba_cuda/numba/cuda/core/rewrites/static_raise.py       | 2 +-
 numba_cuda/numba/cuda/core/ssa.py                         | 2 +-
 numba_cuda/numba/cuda/core/transforms.py                  | 2 +-
 numba_cuda/numba/cuda/core/unsafe/eh.py                   | 2 +-
 numba_cuda/numba/cuda/cpython/listobj.py                  | 2 +-
 numba_cuda/numba/cuda/cpython/numbers.py                  | 2 +-
 numba_cuda/numba/cuda/cpython/unsafe/numbers.py           | 2 +-
 numba_cuda/numba/cuda/cpython/unsafe/tuple.py             | 2 +-
 numba_cuda/numba/cuda/extending.py                        | 2 +-
 numba_cuda/numba/cuda/memory_management/nrt_context.py    | 2 +-
 numba_cuda/numba/cuda/misc/gdb_hook.py                    | 2 +-
 numba_cuda/numba/cuda/np/arrayobj.py                      | 2 +-
 numba_cuda/numba/cuda/np/math/numbers.py                  | 2 +-
 numba_cuda/numba/cuda/np/npyfuncs.py                      | 2 +-
 numba_cuda/numba/cuda/np/npyimpl.py                       | 2 +-
 numba_cuda/numba/cuda/np/numpy_support.py                 | 2 +-
 .../numba/cuda/np/polynomial/polynomial_functions.py      | 2 +-
 numba_cuda/numba/cuda/tests/cudapy/test_ir.py             | 2 +-
 numba_cuda/numba/cuda/tests/cudapy/test_ssa.py            | 2 +-
 numba_cuda/numba/cuda/tests/cudapy/test_typeinfer.py      | 2 +-
 numba_cuda/numba/cuda/tests/support.py                    | 2 +-
 numba_cuda/numba/cuda/tests/test_extending.py             | 8 +++++++-
 numba_cuda/numba/cuda/types/cuda_function_type.py         | 2 +-
 numba_cuda/numba/cuda/types/cuda_functions.py             | 2 +-
 numba_cuda/numba/cuda/typing/asnumbatype.py               | 2 +-
 numba_cuda/numba/cuda/typing/builtins.py                  | 2 +-
 numba_cuda/numba/cuda/typing/collections.py               | 2 +-
 numba_cuda/numba/cuda/typing/dictdecl.py                  | 2 +-
 numba_cuda/numba/cuda/typing/npdatetime.py                | 2 +-
 numba_cuda/numba/cuda/typing/typeof.py                    | 2 +-
 40 files changed, 52 insertions(+), 39 deletions(-)

diff --git a/numba_cuda/numba/cuda/__init__.py b/numba_cuda/numba/cuda/__init__.py
index 04793843c..f1f5335f3 100644
--- a/numba_cuda/numba/cuda/__init__.py
+++ b/numba_cuda/numba/cuda/__init__.py
@@ -71,3 +71,10 @@
 from numba.cuda.np.ufunc import vectorize, guvectorize
 from numba.cuda.misc import quicksort, mergesort
 from numba.cuda.misc.special import literal_unroll
+
+try:
+    import numba
+
+    _HAS_NUMBA = True
+except ImportError:
+    _HAS_NUMBA = False
diff --git a/numba_cuda/numba/cuda/core/analysis.py b/numba_cuda/numba/cuda/core/analysis.py
index d998cc34b..b1bfdaf87 100644
--- a/numba_cuda/numba/cuda/core/analysis.py
+++ b/numba_cuda/numba/cuda/core/analysis.py
@@ -4,7 +4,7 @@
 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
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..f1678caff 100644
--- a/numba_cuda/numba/cuda/core/bytecode.py
+++ b/numba_cuda/numba/cuda/core/bytecode.py
@@ -8,7 +8,7 @@
 
 from types import CodeType, ModuleType
 
-from numba.core import errors
+from numba.cuda.core import errors
 from numba.core import serialize
 from numba.cuda import utils
 from numba.cuda.utils import PYVERSION
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/inline_closurecall.py b/numba_cuda/numba/cuda/core/inline_closurecall.py
index db9f0eb08..66a51b8d5 100644
--- a/numba_cuda/numba/cuda/core/inline_closurecall.py
+++ b/numba_cuda/numba/cuda/core/inline_closurecall.py
@@ -7,7 +7,7 @@
 import numba.cuda.core.analysis
 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,
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 281085e32..b59adbdec 100644
--- a/numba_cuda/numba/cuda/core/ir.py
+++ b/numba_cuda/numba/cuda/core/ir.py
@@ -15,7 +15,7 @@
 from functools import total_ordering
 from io import StringIO
 
-from numba.core import errors
+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 (
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/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/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/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/cpython/listobj.py b/numba_cuda/numba/cuda/cpython/listobj.py
index b1904e514..140206591 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,
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/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/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/arrayobj.py b/numba_cuda/numba/cuda/np/arrayobj.py
index 6df416667..351959ce2 100644
--- a/numba_cuda/numba/cuda/np/arrayobj.py
+++ b/numba_cuda/numba/cuda/np/arrayobj.py
@@ -18,7 +18,7 @@
 
 from numba.cuda 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,
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 6b57f4ef8..d64e9c104 100644
--- a/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py
+++ b/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py
@@ -11,7 +11,7 @@
 
 from numba.cuda 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/tests/cudapy/test_ir.py b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
index 7fac682cb..f0acc1e03 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
@@ -9,7 +9,7 @@
 from numba import objmode
 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,
 )
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/support.py b/numba_cuda/numba/cuda/tests/support.py
index 1a7afeb2b..d9ab68c95 100644
--- a/numba_cuda/numba/cuda/tests/support.py
+++ b/numba_cuda/numba/cuda/tests/support.py
@@ -23,7 +23,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
diff --git a/numba_cuda/numba/cuda/tests/test_extending.py b/numba_cuda/numba/cuda/tests/test_extending.py
index d26a07957..17331693b 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,
 )
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/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 70b7ef532..c380cbf43 100644
--- a/numba_cuda/numba/cuda/typing/builtins.py
+++ b/numba_cuda/numba/cuda/typing/builtins.py
@@ -7,7 +7,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 (
     AttributeTemplate,
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/typeof.py b/numba_cuda/numba/cuda/typing/typeof.py
index f199c5af4..3c718c9d4 100644
--- a/numba_cuda/numba/cuda/typing/typeof.py
+++ b/numba_cuda/numba/cuda/typing/typeof.py
@@ -10,7 +10,7 @@
 from numpy.random.bit_generator import BitGenerator
 
 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
 

From 51372614961b357a496e96bd01ef2a80b463ba28 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Tue, 28 Oct 2025 15:39:31 -0700
Subject: [PATCH 09/25] [NFC][Refactor] Replace some dangling imports, add
 proper guards around numba not being alongside Numba-CUDA

---
 numba_cuda/numba/cuda/__init__.py              | 15 ++++++++-------
 numba_cuda/numba/cuda/compiler.py              |  2 +-
 numba_cuda/numba/cuda/core/analysis.py         |  2 +-
 numba_cuda/numba/cuda/core/byteflow.py         |  2 +-
 numba_cuda/numba/cuda/core/controlflow.py      |  2 +-
 numba_cuda/numba/cuda/core/entrypoints.py      |  6 ++----
 numba_cuda/numba/cuda/core/errors.py           |  2 +-
 numba_cuda/numba/cuda/core/event.py            |  2 +-
 numba_cuda/numba/cuda/core/imputils.py         |  2 +-
 numba_cuda/numba/cuda/core/ir.py               | 14 +++++++++++---
 numba_cuda/numba/cuda/core/tracing.py          |  2 +-
 numba_cuda/numba/cuda/core/typed_passes.py     |  2 +-
 numba_cuda/numba/cuda/cpython/listobj.py       |  2 +-
 numba_cuda/numba/cuda/misc/coverage_support.py |  2 +-
 14 files changed, 32 insertions(+), 25 deletions(-)

diff --git a/numba_cuda/numba/cuda/__init__.py b/numba_cuda/numba/cuda/__init__.py
index f1f5335f3..6301c1b51 100644
--- a/numba_cuda/numba/cuda/__init__.py
+++ b/numba_cuda/numba/cuda/__init__.py
@@ -16,6 +16,14 @@
 # Re-export all type names
 from numba.cuda.types import *
 
+# Check if upstream numba is available - needs to be defined early to avoid circular imports
+try:
+    import numba
+
+    _HAS_NUMBA = True
+except ImportError:
+    _HAS_NUMBA = False
+
 
 # Require NVIDIA CUDA bindings at import time
 if not (
@@ -71,10 +79,3 @@
 from numba.cuda.np.ufunc import vectorize, guvectorize
 from numba.cuda.misc import quicksort, mergesort
 from numba.cuda.misc.special import literal_unroll
-
-try:
-    import numba
-
-    _HAS_NUMBA = True
-except ImportError:
-    _HAS_NUMBA = False
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 b1bfdaf87..bda7fdf2f 100644
--- a/numba_cuda/numba/cuda/core/analysis.py
+++ b/numba_cuda/numba/cuda/core/analysis.py
@@ -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/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/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..6fc0b8377 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 e2303263d..55432141c 100644
--- a/numba_cuda/numba/cuda/core/errors.py
+++ b/numba_cuda/numba/cuda/core/errors.py
@@ -815,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/ir.py b/numba_cuda/numba/cuda/core/ir.py
index b59adbdec..2cdc74961 100644
--- a/numba_cuda/numba/cuda/core/ir.py
+++ b/numba_cuda/numba/cuda/core/ir.py
@@ -25,6 +25,7 @@
     ConstantInferenceError,
 )
 from numba.cuda.core import consts
+from numba.cuda import _HAS_NUMBA
 
 # terminal color markup
 _termcolor = errors.termcolor()
@@ -776,7 +777,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)
         self.value = value
         self.loc = loc
 
@@ -1342,8 +1346,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
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/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/cpython/listobj.py b/numba_cuda/numba/cuda/cpython/listobj.py
index 140206591..de28b6399 100644
--- a/numba_cuda/numba/cuda/cpython/listobj.py
+++ b/numba_cuda/numba/cuda/cpython/listobj.py
@@ -21,7 +21,7 @@
 from numba.cuda.extending import overload_method, overload
 from numba.cuda.misc import quicksort
 from numba.cuda.cpython import slicing
-from numba import literal_unroll
+from numba.cuda import literal_unroll
 
 registry = Registry("listobj")
 lower = registry.lower
diff --git a/numba_cuda/numba/cuda/misc/coverage_support.py b/numba_cuda/numba/cuda/misc/coverage_support.py
index 1167fd696..5b50c9c85 100644
--- a/numba_cuda/numba/cuda/misc/coverage_support.py
+++ b/numba_cuda/numba/cuda/misc/coverage_support.py
@@ -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:

From 3c85fde9d32cd0e434cc8bdd6f752be44b5b452f Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Tue, 28 Oct 2025 16:45:19 -0700
Subject: [PATCH 10/25] [Refactor][NFC] Introduce version_info structure and
 version generation function; update imports to use numba.cuda (berzerk
 cleaning of imports, introduced # compat-ignore as a way for my scripts to
 ignore necessary numba imports)

---
 numba_cuda/_version.py                        | 44 +++++++++++++++++++
 numba_cuda/numba/cuda/__init__.py             |  7 +++
 numba_cuda/numba/cuda/core/bytecode.py        |  2 +-
 numba_cuda/numba/cuda/core/config.py          |  2 +-
 numba_cuda/numba/cuda/core/entrypoints.py     |  2 +-
 numba_cuda/numba/cuda/core/errors.py          |  4 +-
 numba_cuda/numba/cuda/core/sigutils.py        | 14 +++---
 numba_cuda/numba/cuda/extending.py            | 28 ++++++------
 numba_cuda/numba/cuda/lowering.py             |  2 +-
 numba_cuda/numba/cuda/misc/quicksort.py       |  2 +-
 numba_cuda/numba/cuda/random.py               |  6 +--
 numba_cuda/numba/cuda/target.py               |  2 +-
 numba_cuda/numba/cuda/testing.py              |  6 +++
 .../numba/cuda/tests/core/test_serialize.py   | 13 +++++-
 .../cuda/tests/cudapy/test_array_alignment.py |  9 +++-
 .../tests/cudapy/test_cuda_array_interface.py |  2 +-
 .../numba/cuda/tests/cudapy/test_datetime.py  |  3 +-
 .../numba/cuda/tests/cudapy/test_enums.py     | 17 ++++++-
 .../numba/cuda/tests/cudapy/test_gufunc.py    |  4 +-
 .../cuda/tests/cudapy/test_gufunc_scalar.py   |  3 +-
 .../cuda/tests/cudapy/test_intrinsics.py      |  7 ++-
 numba_cuda/numba/cuda/tests/cudapy/test_ir.py | 23 ++++++----
 .../numba/cuda/tests/cudapy/test_localmem.py  |  7 ++-
 .../numba/cuda/tests/cudapy/test_math.py      |  3 +-
 .../numba/cuda/tests/cudapy/test_overload.py  | 25 +++++++++--
 .../numba/cuda/tests/cudapy/test_serialize.py |  3 +-
 numba_cuda/numba/cuda/tests/cudapy/test_sm.py | 13 ++++--
 .../cuda/tests/cudapy/test_sm_creation.py     |  7 ++-
 .../numba/cuda/tests/cudapy/test_vectorize.py |  2 +-
 numba_cuda/numba/cuda/tests/support.py        | 12 +++--
 .../numba/cuda/tests/test_extending_types.py  |  7 ++-
 numba_cuda/numba/cuda/typing/templates.py     | 12 +++--
 32 files changed, 214 insertions(+), 79 deletions(-)

diff --git a/numba_cuda/_version.py b/numba_cuda/_version.py
index b5622732a..83d5a3894 100644
--- a/numba_cuda/_version.py
+++ b/numba_cuda/_version.py
@@ -2,6 +2,7 @@
 # SPDX-License-Identifier: BSD-2-Clause
 
 import importlib.resources
+from collections import namedtuple
 
 __version__ = (
     importlib.resources.files("numba_cuda")
@@ -9,3 +10,46 @@
     .read_text()
     .strip()
 )
+
+version_info = namedtuple(
+    "version_info", ("major minor patch short full string tuple git_revision")
+)
+
+
+def generate_version_info(version):
+    """Process a version string into a structured version_info object.
+
+    Parameters
+    ----------
+    version: str
+        a string describing the current version
+
+    Returns
+    -------
+    version_info: tuple
+        structured version information
+
+    See also
+    --------
+    Look at the definition of 'version_info' in this module for details.
+
+    """
+    parts = version.split(".")
+
+    def try_int(x):
+        try:
+            return int(x)
+        except ValueError:
+            return None
+
+    major = try_int(parts[0]) if len(parts) >= 1 else None
+    minor = try_int(parts[1]) if len(parts) >= 2 else None
+    patch = try_int(parts[2]) if len(parts) >= 3 else None
+    short = (major, minor)
+    full = (major, minor, patch)
+    string = version
+    tup = tuple(parts)
+    git_revision = tup[3] if len(tup) >= 4 else None
+    return version_info(
+        major, minor, patch, short, full, string, tup, git_revision
+    )
diff --git a/numba_cuda/numba/cuda/__init__.py b/numba_cuda/numba/cuda/__init__.py
index 6301c1b51..f7ed99364 100644
--- a/numba_cuda/numba/cuda/__init__.py
+++ b/numba_cuda/numba/cuda/__init__.py
@@ -10,6 +10,12 @@
 # Import version from the parent package
 from numba_cuda._version import __version__
 
+# Generate version_info early to avoid circular import issues
+from numba_cuda._version import generate_version_info
+
+version_info = generate_version_info(__version__)
+del generate_version_info
+
 # Re-export types itself
 import numba.cuda.types as types
 
@@ -79,3 +85,4 @@
 from numba.cuda.np.ufunc import vectorize, guvectorize
 from numba.cuda.misc import quicksort, mergesort
 from numba.cuda.misc.special import literal_unroll
+from numba.cuda.np.numpy_support import *
diff --git a/numba_cuda/numba/cuda/core/bytecode.py b/numba_cuda/numba/cuda/core/bytecode.py
index f1678caff..d86b6b496 100644
--- a/numba_cuda/numba/cuda/core/bytecode.py
+++ b/numba_cuda/numba/cuda/core/bytecode.py
@@ -9,7 +9,7 @@
 from types import CodeType, ModuleType
 
 from numba.cuda.core import errors
-from numba.core import serialize
+from numba.cuda import serialize
 from numba.cuda import utils
 from numba.cuda.utils import PYVERSION
 
diff --git a/numba_cuda/numba/cuda/core/config.py b/numba_cuda/numba/cuda/core/config.py
index 92db61688..8f97a90dc 100644
--- a/numba_cuda/numba/cuda/core/config.py
+++ b/numba_cuda/numba/cuda/core/config.py
@@ -648,7 +648,7 @@ def reload_config():
 
 # use numba.core.config if available, otherwise use numba.cuda.core.config
 try:
-    import numba.core.config as _config
+    import numba.core.config as _config  # compat-ignore
 
     sys.modules[__name__] = _config
 except ImportError:
diff --git a/numba_cuda/numba/cuda/core/entrypoints.py b/numba_cuda/numba/cuda/core/entrypoints.py
index 6fc0b8377..279195ca5 100644
--- a/numba_cuda/numba/cuda/core/entrypoints.py
+++ b/numba_cuda/numba/cuda/core/entrypoints.py
@@ -17,7 +17,7 @@ def init_all():
     If extensions have already been initialized, this function does nothing.
     """
     if _HAS_NUMBA:
-        from numba.core import entrypoints
+        from numba.core import entrypoints  # compat-ignore
 
         entrypoints.init_all()
 
diff --git a/numba_cuda/numba/cuda/core/errors.py b/numba_cuda/numba/cuda/core/errors.py
index 55432141c..68c63f851 100644
--- a/numba_cuda/numba/cuda/core/errors.py
+++ b/numba_cuda/numba/cuda/core/errors.py
@@ -9,7 +9,7 @@
 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
@@ -26,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:
diff --git a/numba_cuda/numba/cuda/core/sigutils.py b/numba_cuda/numba/cuda/core/sigutils.py
index 04fa54c45..897275040 100644
--- a/numba_cuda/numba/cuda/core/sigutils.py
+++ b/numba_cuda/numba/cuda/core/sigutils.py
@@ -1,14 +1,10 @@
 # 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:
-    from numba.core.typing import Signature as CoreSignature
-
-    numba_sig_present = True
-except ImportError:
-    numba_sig_present = False
+if _HAS_NUMBA:
+    from numba.core.typing import Signature as CoreSignature  # compat-ignore
 
 
 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/extending.py b/numba_cuda/numba/cuda/extending.py
index 11de8d096..9ce1f6ad5 100644
--- a/numba_cuda/numba/cuda/extending.py
+++ b/numba_cuda/numba/cuda/extending.py
@@ -219,12 +219,12 @@ def decorate(overload_func):
         # For generic/CPU targets, also register in numba.core registry
         if target in ("generic", "cpu"):
             try:
-                from numba.core.typing.templates import (
-                    make_overload_template as core_make_overload_template,
-                    infer as core_infer,
-                    infer_global as core_infer_global,
-                )
-                from numba.core import types as core_types
+                from numba.core.typing.templates import (  # compat-ignore
+                    make_overload_template as core_make_overload_template,  # compat-ignore
+                    infer as core_infer,  # compat-ignore
+                    infer_global as core_infer_global,  # compat-ignore
+                )  # compat-ignore
+                from numba.core import types as core_types  # compat-ignore
 
                 core_template = core_make_overload_template(
                     func,
@@ -310,10 +310,10 @@ def decorate(overload_func):
         # For generic/CPU targets, also register in numba.core registry
         if target in ("generic", "cpu"):
             try:
-                from numba.core.typing.templates import (
-                    make_overload_attribute_template as core_make_overload_attribute_template,
-                    infer_getattr as core_infer_getattr,
-                )
+                from numba.core.typing.templates import (  # compat-ignore
+                    make_overload_attribute_template as core_make_overload_attribute_template,  # compat-ignore
+                    infer_getattr as core_infer_getattr,  # compat-ignore
+                )  # compat-ignore
 
                 core_template = core_make_overload_attribute_template(
                     typ, attr, overload_func, **kwargs
@@ -348,10 +348,10 @@ def decorate(overload_func):
         target = kwargs.get("target", "cuda")
         if target in ("generic", "cpu"):
             try:
-                from numba.core.typing.templates import (
-                    make_overload_method_template as core_make_overload_method_template,
-                    infer_getattr as core_infer_getattr,
-                )
+                from numba.core.typing.templates import (  # compat-ignore
+                    make_overload_method_template as core_make_overload_method_template,  # compat-ignore
+                    infer_getattr as core_infer_getattr,  # compat-ignore
+                )  # compat-ignore
 
                 copied_kwargs2 = kwargs.copy()
                 core_template = core_make_overload_method_template(
diff --git a/numba_cuda/numba/cuda/lowering.py b/numba_cuda/numba/cuda/lowering.py
index 16a557dbc..2bf11362a 100644
--- a/numba_cuda/numba/cuda/lowering.py
+++ b/numba_cuda/numba/cuda/lowering.py
@@ -32,7 +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
+from numba.cuda import version_info
 
 numba_version = version_info.short
 del version_info
diff --git a/numba_cuda/numba/cuda/misc/quicksort.py b/numba_cuda/numba/cuda/misc/quicksort.py
index 042fc2aba..888ace810 100644
--- a/numba_cuda/numba/cuda/misc/quicksort.py
+++ b/numba_cuda/numba/cuda/misc/quicksort.py
@@ -268,6 +268,6 @@ def make_py_quicksort(*args, **kwargs):
 
 
 def make_jit_quicksort(*args, **kwargs):
-    from numba.core.extending import register_jitable
+    from numba.cuda.extending import register_jitable
 
     return make_quicksort_impl((lambda f: register_jitable(f)), *args, **kwargs)
diff --git a/numba_cuda/numba/cuda/random.py b/numba_cuda/numba/cuda/random.py
index 47a76b248..dcfc6bff4 100644
--- a/numba_cuda/numba/cuda/random.py
+++ b/numba_cuda/numba/cuda/random.py
@@ -3,17 +3,17 @@
 
 import math
 
-from numba import (
-    cuda,
+import numba.cuda as cuda
+from numba.cuda import (
     float32,
     float64,
     uint32,
     int64,
     uint64,
     from_dtype,
-    jit,
 )
 from numba.cuda import config
+from numba import jit  # compat-ignore
 
 import numpy as np
 
diff --git a/numba_cuda/numba/cuda/target.py b/numba_cuda/numba/cuda/target.py
index 6faf15571..ac3d875b5 100644
--- a/numba_cuda/numba/cuda/target.py
+++ b/numba_cuda/numba/cuda/target.py
@@ -226,7 +226,7 @@ def load_additional_registries(self):
         # Install only implementations that are defined outside of numba (i.e.,
         # in third-party extensions) from Numba's builtin_registry.
         if importlib.util.find_spec("numba.core.imputils") is not None:
-            from numba.core.imputils import builtin_registry
+            from numba.core.imputils import builtin_registry  # compat-ignore
 
             self.install_external_registry(builtin_registry)
 
diff --git a/numba_cuda/numba/cuda/testing.py b/numba_cuda/numba/cuda/testing.py
index 196d08897..caf93a831 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/test_serialize.py b/numba_cuda/numba/cuda/tests/core/test_serialize.py
index f5db16869..b5a49d2d6 100644
--- a/numba_cuda/numba/cuda/tests/core/test_serialize.py
+++ b/numba_cuda/numba/cuda/tests/core/test_serialize.py
@@ -10,10 +10,16 @@
 import unittest
 from multiprocessing import get_context
 
-import numba
-from numba.core.errors import TypingError
+from numba.cuda import _HAS_NUMBA
+
+if _HAS_NUMBA:
+    from numba.core.errors import TypingError  # compat-ignore
+    import numba  # compat-ignore
+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
 
 try:
     from numba.core.target_extension import resolve_dispatcher_from_str
@@ -304,7 +310,10 @@ class Klass:
         proc.join(timeout=60)
         self.assertEqual(proc.exitcode, 0)
 
+    @skip_on_standalone_numba_cuda
     def test_dynamic_class_issue_7356(self):
+        import numba  # compat-ignore
+
         cfunc = numba.njit(issue_7356)
         self.assertEqual(cfunc(), (100, 100))
 
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..e90ed929d 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_array_alignment.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_array_alignment.py
@@ -4,8 +4,13 @@
 import re
 import itertools
 import numpy as np
-from numba import cuda
-from numba.core.errors import TypingError
+import numba.cuda as cuda
+from numba.cuda import _HAS_NUMBA
+
+if _HAS_NUMBA:
+    from numba.core.errors import TypingError  # compat-ignore
+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_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_datetime.py b/numba_cuda/numba/cuda/tests/cudapy/test_datetime.py
index 2c513dd0a..e44522416 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_datetime.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_datetime.py
@@ -3,7 +3,8 @@
 
 import numpy as np
 
-from numba import cuda, vectorize, guvectorize
+from numba.cuda import vectorize, guvectorize
+import numba.cuda as cuda
 from numba.cuda.np.numpy_support import from_dtype
 from numba.cuda.testing import CUDATestCase, skip_on_cudasim
 import unittest
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_enums.py b/numba_cuda/numba/cuda/tests/cudapy/test_enums.py
index 12c5161b4..e1e035c3c 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_enums.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_enums.py
@@ -8,9 +8,19 @@
 import numpy as np
 
 from numba.cuda import int16, int32
-from numba import cuda, vectorize, njit
+from numba.cuda import vectorize
+from numba.cuda import _HAS_NUMBA
+
+if _HAS_NUMBA:
+    from numba import njit  # compat-ignore
+import numba.cuda as 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 +67,9 @@ def f(out):
         f(expected)
         self.assertPreciseEqual(expected, got)
 
+    @skip_on_standalone_numba_cuda(
+        "Test not supported in standalone numba_cuda"
+    )
     def test_return_from_device_func(self):
         @njit
         def inner(pred):
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_gufunc.py b/numba_cuda/numba/cuda/tests/cudapy/test_gufunc.py
index 0d7b1c075..019482321 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_gufunc.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_gufunc.py
@@ -5,8 +5,8 @@
 
 from collections import namedtuple
 from numba.cuda import void, int32, float32, float64
-from numba import guvectorize
-from numba import cuda
+from numba.cuda import guvectorize
+import numba.cuda as cuda
 from numba.cuda.testing import skip_on_cudasim, CUDATestCase
 import unittest
 import warnings
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..7cb4ddca3 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,8 @@
 """
 
 import numpy as np
-from numba import guvectorize, cuda
+from numba.cuda import guvectorize
+import numba.cuda as cuda
 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..d8490cefe 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  # compat-ignore
+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 f0acc1e03..e4bc24b26 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
@@ -6,7 +6,11 @@
 import warnings
 import numpy as np
 
-from numba import objmode
+from numba import _HAS_NUMBA
+
+if _HAS_NUMBA:
+    from numba import objmode  # compat-ignore
+    from numba import njit  # compat-ignore
 from numba.cuda.core import ir
 from numba.cuda import compiler
 from numba.cuda.core import errors
@@ -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,6 @@ def gen_block():
 
         self.check(a, same=[b], different=[c])
 
-    @requires_fully_vendored_ir
     def test_functionir(self):
         def run_frontend(x):
             return compiler.run_frontend(x, emit_dels=True)
@@ -396,8 +399,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 +415,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 +519,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..4465223e3 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  # compat-ignore
+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_math.py b/numba_cuda/numba/cuda/tests/cudapy/test_math.py
index 71b089cf4..97c2a9613 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_math.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_math.py
@@ -10,7 +10,8 @@
     skip_on_cudasim,
 )
 from numba.cuda.np import numpy_support
-from numba import cuda, vectorize
+from numba import vectorize
+import numba.cuda as cuda
 from numba.cuda import float32, float64, int32, void, int64
 import math
 
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_overload.py b/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
index ff8190671..7176dc7b3 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
+import numba.cuda as cuda
+import numba.cuda.types as types
+from numba.cuda import _HAS_NUMBA
+
+if _HAS_NUMBA:
+    from numba.core.errors import TypingError  # compat-ignore
+    from numba import njit  # compat-ignore
+else:
+    from numba.cuda.core.errors import TypingError
+from numba.cuda import version_info
 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,9 @@ def check_overload(self, kernel, expected):
         cuda.jit(kernel)[1, 1](x)
         self.assertEqual(x[0], expected)
 
+    @skip_on_standalone_numba_cuda(
+        "Test not supported in standalone numba_cuda"
+    )
     def check_overload_cpu(self, kernel, expected):
         x = np.ones(1, dtype=np.int32)
         njit(kernel)(x)
@@ -328,6 +344,9 @@ def kernel(x):
         expected = GENERIC_TARGET_OL_CALLS_TARGET_OL * GENERIC_TARGET_OL
         self.check_overload_cpu(kernel, expected)
 
+    @skip_on_standalone_numba_cuda(
+        "Test not supported in standalone numba_cuda"
+    )
     def test_overload_attribute_target(self):
         MyDummy, MyDummyType = self.make_dummy_type()
         mydummy_type_cpu = cpu_typeof(MyDummy())  # For @njit (cpu)
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_serialize.py b/numba_cuda/numba/cuda/tests/cudapy/test_serialize.py
index f82abfc67..0429b0013 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_serialize.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_serialize.py
@@ -3,7 +3,8 @@
 
 import pickle
 import numpy as np
-from numba import cuda, vectorize
+from numba.cuda import vectorize
+import numba.cuda as cuda
 from numba.cuda import types
 from numba.cuda.testing import skip_on_cudasim, CUDATestCase
 import unittest
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_sm.py b/numba_cuda/numba/cuda/tests/cudapy/test_sm.py
index 85424fcdd..a58922da6 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_sm.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_sm.py
@@ -3,8 +3,13 @@
 
 from numba import cuda
 from numba.cuda import int32, float64, void
-from numba.cuda.core.errors import TypingError
-import numba.core.errors as numbaError
+from numba.cuda import _HAS_NUMBA
+
+if _HAS_NUMBA:
+    from numba.core.errors import TypingError  # compat-ignore
+    from numba.cuda.core.errors import TypingError as cudaTypingError
+else:
+    from numba.cuda.core.errors import TypingError as cudaTypingError
 from numba.cuda import types
 from numba.cuda.testing import unittest, CUDATestCase, skip_on_cudasim
 
@@ -402,7 +407,7 @@ def test_invalid_array_type(self):
         def unsupported_type():
             arr = cuda.shared.array(10, dtype=np.dtype("O"))  # noqa: F841
 
-        with self.assertRaisesRegex(TypingError, rgx):
+        with self.assertRaisesRegex(cudaTypingError, rgx):
             cuda.jit(void())(unsupported_type)
 
         rgx = ".*Invalid NumPy dtype specified: 'int33'.*"
@@ -410,7 +415,7 @@ def unsupported_type():
         def invalid_string_type():
             arr = cuda.shared.array(10, dtype="int33")  # noqa: F841
 
-        with self.assertRaisesRegex(numbaError.TypingError, rgx):
+        with self.assertRaisesRegex(TypingError, rgx):
             cuda.jit(void())(invalid_string_type)
 
     @skip_on_cudasim("Struct model array unsupported in simulator")
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..6b3d156d5 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  # compat-ignore
+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_vectorize.py b/numba_cuda/numba/cuda/tests/cudapy/test_vectorize.py
index 3a41b1234..9b779e248 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_vectorize.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_vectorize.py
@@ -7,7 +7,7 @@
 from functools import partial
 from itertools import product
 from numba.cuda import vectorize as cuda_vectorize
-from numba import cuda, vectorize as numba_vectorize
+from numba import cuda, vectorize as numba_vectorize  # compat-ignore
 from numba.cuda.types import int32, float32, float64
 from numba.cuda.cudadrv.driver import CudaAPIError, driver
 from numba.cuda.testing import skip_on_cudasim
diff --git a/numba_cuda/numba/cuda/tests/support.py b/numba_cuda/numba/cuda/tests/support.py
index d9ab68c95..b9ad46c18 100644
--- a/numba_cuda/numba/cuda/tests/support.py
+++ b/numba_cuda/numba/cuda/tests/support.py
@@ -36,10 +36,14 @@
 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.core import types as upstream_types
-except ImportError:
+from numba.cuda import _HAS_NUMBA
+
+if _HAS_NUMBA:
+    from numba.core.extending import (
+        typeof_impl as upstream_typeof_impl,
+    )  # compat-ignore
+    from numba.core import types as upstream_types  # compat-ignore
+else:
     upstream_typeof_impl = None
     upstream_types = None
 
diff --git a/numba_cuda/numba/cuda/tests/test_extending_types.py b/numba_cuda/numba/cuda/tests/test_extending_types.py
index 08dfd4da2..0c78efddf 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  # compat-ignore
+else:
+    from numba.cuda.core.errors import TypingError
 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/typing/templates.py b/numba_cuda/numba/cuda/typing/templates.py
index 235bc758f..3f78298c6 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  # compat-ignore
 
 # 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(

From 37692da353b4530598f18b39232b70602c3c90bb Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Wed, 29 Oct 2025 10:32:14 -0700
Subject: [PATCH 11/25] update numba.cuda imports to fix test errors

---
 numba_cuda/numba/cuda/tests/cudapy/test_ipc.py      | 2 +-
 numba_cuda/numba/cuda/tests/cudapy/test_ir.py       | 2 +-
 numba_cuda/numba/cuda/tests/cudapy/test_overload.py | 6 +++---
 numba_cuda/numba/cuda/tests/cudapy/test_warning.py  | 2 +-
 numba_cuda/numba/cuda/tests/nrt/test_nrt.py         | 6 +++---
 5 files changed, 9 insertions(+), 9 deletions(-)

diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_ipc.py b/numba_cuda/numba/cuda/tests/cudapy/test_ipc.py
index 3f21dd7bd..4f27b3a88 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_ipc.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_ipc.py
@@ -10,7 +10,7 @@
 
 import numpy as np
 
-from numba import cuda
+import numba.cuda as cuda
 from numba.cuda.testing import (
     skip_on_arm,
     skip_on_cudasim,
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_ir.py b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
index e4bc24b26..ea7da7c6e 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
@@ -6,7 +6,7 @@
 import warnings
 import numpy as np
 
-from numba import _HAS_NUMBA
+from numba.cuda import _HAS_NUMBA
 
 if _HAS_NUMBA:
     from numba import objmode  # compat-ignore
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_overload.py b/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
index 7176dc7b3..ff2e3a89b 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
@@ -8,12 +8,12 @@
 if _HAS_NUMBA:
     from numba.core.errors import TypingError  # compat-ignore
     from numba import njit  # compat-ignore
+    import numba  # compat-ignore
 else:
     from numba.cuda.core.errors import TypingError
-from numba.cuda import version_info
 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.core.typing.typeof import typeof as cpu_typeof  # compat-ignore
 from numba.cuda.testing import (
     CUDATestCase,
     skip_on_cudasim,
@@ -365,7 +365,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_warning.py b/numba_cuda/numba/cuda/tests/cudapy/test_warning.py
index a80278eaf..3d2b1eb3e 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_warning.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_warning.py
@@ -2,7 +2,7 @@
 # SPDX-License-Identifier: BSD-2-Clause
 
 import numpy as np
-from numba import cuda
+import numba.cuda as cuda
 from numba.cuda.cudadrv import driver
 from numba.cuda.testing import (
     unittest,
diff --git a/numba_cuda/numba/cuda/tests/nrt/test_nrt.py b/numba_cuda/numba/cuda/tests/nrt/test_nrt.py
index 1efeed919..b4a45d7a1 100644
--- a/numba_cuda/numba/cuda/tests/nrt/test_nrt.py
+++ b/numba_cuda/numba/cuda/tests/nrt/test_nrt.py
@@ -12,7 +12,7 @@
 from numba.cuda.cudadrv.nvrtc import compile
 from numba.cuda import types
 from numba.cuda.typing import signature
-from numba import cuda
+import numba.cuda as cuda
 from numba.cuda import config
 from numba.cuda.typing.templates import AbstractTemplate
 from numba.cuda.cudadrv.linkable_code import (
@@ -218,7 +218,7 @@ def tearDown(self):
     def test_stats_env_var_explicit_on(self):
         # Checks that explicitly turning the stats on via the env var works.
         src = """if 1:
-        from numba import cuda
+        import numba.cuda as cuda
         from numba.cuda.memory_management import rtsys
         import numpy as np
 
@@ -256,7 +256,7 @@ def foo():
 
     def check_env_var_off(self, env):
         src = """if 1:
-        from numba import cuda
+        import numba.cuda as cuda
         import numpy as np
         from numba.cuda.memory_management import rtsys
 

From c7239fd4008454beb230c966c8c8c9d52214e8d5 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Wed, 29 Oct 2025 11:07:36 -0700
Subject: [PATCH 12/25] remove legacy_type_system usage

---
 numba_cuda/numba/cuda/misc/quicksort.py | 7 ++-----
 1 file changed, 2 insertions(+), 5 deletions(-)

diff --git a/numba_cuda/numba/cuda/misc/quicksort.py b/numba_cuda/numba/cuda/misc/quicksort.py
index 888ace810..506e3636d 100644
--- a/numba_cuda/numba/cuda/misc/quicksort.py
+++ b/numba_cuda/numba/cuda/misc/quicksort.py
@@ -5,7 +5,7 @@
 
 import numpy as np
 
-from numba.cuda import types, config
+from numba.cuda import types
 
 
 QuicksortImplementation = collections.namedtuple(
@@ -33,10 +33,7 @@
 def make_quicksort_impl(
     wrap, lt=None, is_argsort=False, is_list=False, is_np_array=False
 ):
-    if config.USE_LEGACY_TYPE_SYSTEM:
-        intp = types.intp
-    else:
-        intp = types.py_int
+    intp = types.intp
     zero = intp(0)
 
     # Two subroutines to make the core algorithm generic wrt. argsort

From c28b06348a187347a586da765a18268df63f08bf Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Wed, 29 Oct 2025 11:08:54 -0700
Subject: [PATCH 13/25] remove LEGACY_TYPE_SYSTEM in numba.cuda.core.config,
 skip test on sim

---
 numba_cuda/numba/cuda/core/config.py          | 5 -----
 numba_cuda/numba/cuda/tests/cudapy/test_ir.py | 3 ++-
 2 files changed, 2 insertions(+), 6 deletions(-)

diff --git a/numba_cuda/numba/cuda/core/config.py b/numba_cuda/numba/cuda/core/config.py
index 8f97a90dc..15fca7f13 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/tests/cudapy/test_ir.py b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
index ea7da7c6e..f7ff4195f 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
@@ -2,7 +2,7 @@
 # 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
 
@@ -378,6 +378,7 @@ def gen_block():
 
         self.check(a, same=[b], different=[c])
 
+    @skip_on_cudasim
     def test_functionir(self):
         def run_frontend(x):
             return compiler.run_frontend(x, emit_dels=True)

From 0a60d27f5a0492c30b50393e3fc1fc395b308b35 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Wed, 29 Oct 2025 13:19:26 -0700
Subject: [PATCH 14/25] Remove MakeFunctionToJitFunction pass, which implicitly
 uses njit, not a use case we intend to support

---
 numba_cuda/numba/cuda/compiler.py            |  6 --
 numba_cuda/numba/cuda/core/untyped_passes.py | 67 --------------------
 2 files changed, 73 deletions(-)

diff --git a/numba_cuda/numba/cuda/compiler.py b/numba_cuda/numba/cuda/compiler.py
index a1653c192..8d76bc675 100644
--- a/numba_cuda/numba/cuda/compiler.py
+++ b/numba_cuda/numba/cuda/compiler.py
@@ -48,7 +48,6 @@
     WithLifting,
     InlineInlinables,
     FindLiterallyCalls,
-    MakeFunctionToJitFunction,
     LiteralUnroll,
     ReconstructSSA,
     RewriteDynamicRaises,
@@ -225,11 +224,6 @@ def define_untyped_pipeline(state, name="untyped"):
 
         pm.add_pass(RewriteDynamicRaises, "rewrite dynamic raises")
 
-        # convert any remaining closures into functions
-        pm.add_pass(
-            MakeFunctionToJitFunction,
-            "convert make_function into JIT functions",
-        )
         # inline functions that have been determined as inlinable and rerun
         # branch pruning, this needs to be run after closures are inlined as
         # the IR repr of a closure masks call sites if an inlinable is called
diff --git a/numba_cuda/numba/cuda/core/untyped_passes.py b/numba_cuda/numba/cuda/core/untyped_passes.py
index 4006854ce..a6f20a2f1 100644
--- a/numba_cuda/numba/cuda/core/untyped_passes.py
+++ b/numba_cuda/numba/cuda/core/untyped_passes.py
@@ -35,7 +35,6 @@
     resolve_func_from_module,
     simplify_CFG,
     GuardException,
-    convert_code_obj_to_function,
     build_definitions,
     replace_var_names,
     get_name_var_table,
@@ -618,72 +617,6 @@ def run_pass(self, state):
         return False
 
 
-@register_pass(mutates_CFG=True, analysis_only=False)
-class MakeFunctionToJitFunction(FunctionPass):
-    """
-    This swaps an ir.Expr.op == "make_function" i.e. a closure, for a compiled
-    function containing the closure body and puts it in ir.Global. It's a 1:1
-    statement value swap. `make_function` is already untyped
-    """
-
-    _name = "make_function_op_code_to_jit_function"
-
-    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():
-            for stmt in blk.body:
-                if isinstance(stmt, ir.Assign):
-                    if isinstance(stmt.value, ir.Expr):
-                        if stmt.value.op == "make_function":
-                            node = stmt.value
-                            getdef = func_ir.get_definition
-                            kw_default = getdef(node.defaults)
-                            ok = False
-                            if kw_default is None or isinstance(
-                                kw_default, ir.Const
-                            ):
-                                ok = True
-                            elif isinstance(kw_default, tuple):
-                                ok = all(
-                                    [
-                                        isinstance(getdef(x), ir.Const)
-                                        for x in kw_default
-                                    ]
-                                )
-                            elif isinstance(kw_default, ir.Expr):
-                                if kw_default.op != "build_tuple":
-                                    continue
-                                ok = all(
-                                    [
-                                        isinstance(getdef(x), ir.Const)
-                                        for x in kw_default.items
-                                    ]
-                                )
-                            if not ok:
-                                continue
-
-                            pyfunc = convert_code_obj_to_function(node, func_ir)
-                            func = njit()(pyfunc)
-                            new_node = ir.Global(
-                                node.code.co_name, func, stmt.loc
-                            )
-                            stmt.value = new_node
-                            mutated |= True
-
-        # if a change was made the del ordering is probably wrong, patch up
-        if mutated:
-            post_proc = postproc.PostProcessor(func_ir)
-            post_proc.run()
-
-        return mutated
-
-
 @register_pass(mutates_CFG=True, analysis_only=False)
 class TransformLiteralUnrollConstListToTuple(FunctionPass):
     """This pass spots a `literal_unroll([<constant values>])` and rewrites it

From 5179b84e4a4f3271c3bd6271945d3e3f073a0ca0 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Wed, 29 Oct 2025 13:55:58 -0700
Subject: [PATCH 15/25] bunch of sporadic import fixes, cleanup soon

---
 numba_cuda/numba/cuda/core/config.py          |  2 +-
 numba_cuda/numba/cuda/core/entrypoints.py     |  2 +-
 .../numba/cuda/core/inline_closurecall.py     | 16 ++--------
 numba_cuda/numba/cuda/core/ir.py              |  5 ++-
 numba_cuda/numba/cuda/core/ir_utils.py        |  5 ++-
 numba_cuda/numba/cuda/core/sigutils.py        |  2 +-
 numba_cuda/numba/cuda/core/targetconfig.py    |  5 +--
 numba_cuda/numba/cuda/extending.py            | 28 ++++++++--------
 numba_cuda/numba/cuda/lowering.py             | 11 +++++--
 numba_cuda/numba/cuda/misc/mergesort.py       | 18 ++---------
 numba_cuda/numba/cuda/np/ufunc/decorators.py  |  2 +-
 numba_cuda/numba/cuda/random.py               |  3 +-
 numba_cuda/numba/cuda/target.py               | 10 ++----
 .../cuda/tests/core/serialize_usecases.py     |  5 ++-
 .../numba/cuda/tests/core/test_serialize.py   | 32 ++++++++++++-------
 .../tests/cudapy/cache_with_cpu_usecases.py   | 13 +++++---
 .../cuda/tests/cudapy/test_array_alignment.py |  4 +--
 .../tests/cudapy/test_bfloat16_bindings.py    |  2 +-
 .../numba/cuda/tests/cudapy/test_caching.py   |  3 ++
 .../numba/cuda/tests/cudapy/test_datetime.py  |  2 +-
 .../cuda/tests/cudapy/test_debuginfo_types.py |  2 +-
 .../cuda/tests/cudapy/test_device_func.py     | 10 +++++-
 .../numba/cuda/tests/cudapy/test_enums.py     |  4 +--
 .../numba/cuda/tests/cudapy/test_extending.py | 12 ++++++-
 .../numba/cuda/tests/cudapy/test_gufunc.py    |  2 +-
 .../cuda/tests/cudapy/test_gufunc_scalar.py   |  2 +-
 .../cuda/tests/cudapy/test_intrinsics.py      |  2 +-
 .../numba/cuda/tests/cudapy/test_ipc.py       |  2 +-
 numba_cuda/numba/cuda/tests/cudapy/test_ir.py |  4 +--
 .../numba/cuda/tests/cudapy/test_localmem.py  |  2 +-
 .../numba/cuda/tests/cudapy/test_overload.py  | 10 +++---
 .../numba/cuda/tests/cudapy/test_serialize.py |  2 +-
 numba_cuda/numba/cuda/tests/cudapy/test_sm.py |  2 +-
 .../cuda/tests/cudapy/test_sm_creation.py     |  2 +-
 .../numba/cuda/tests/cudapy/test_ufuncs.py    |  7 +++-
 .../numba/cuda/tests/cudapy/test_vectorize.py |  2 +-
 .../numba/cuda/tests/cudapy/test_warning.py   |  2 +-
 .../tests/doc_examples/test_cpu_gpu_compat.py |  7 +++-
 .../tests/doc_examples/test_montecarlo.py     |  7 +++-
 .../numba/cuda/tests/nocuda/test_import.py    |  2 --
 numba_cuda/numba/cuda/tests/nrt/test_nrt.py   |  6 ++--
 numba_cuda/numba/cuda/tests/support.py        |  4 +--
 .../numba/cuda/tests/test_array_reductions.py |  2 +-
 .../numba/cuda/tests/test_extending_types.py  |  2 +-
 numba_cuda/numba/cuda/typing/templates.py     |  2 +-
 numba_cuda/numba/cuda/typing/typeof.py        | 16 ++--------
 46 files changed, 153 insertions(+), 134 deletions(-)

diff --git a/numba_cuda/numba/cuda/core/config.py b/numba_cuda/numba/cuda/core/config.py
index 15fca7f13..5774c8b79 100644
--- a/numba_cuda/numba/cuda/core/config.py
+++ b/numba_cuda/numba/cuda/core/config.py
@@ -643,7 +643,7 @@ def reload_config():
 
 # use numba.core.config if available, otherwise use numba.cuda.core.config
 try:
-    import numba.core.config as _config  # compat-ignore
+    import numba.core.config as _config
 
     sys.modules[__name__] = _config
 except ImportError:
diff --git a/numba_cuda/numba/cuda/core/entrypoints.py b/numba_cuda/numba/cuda/core/entrypoints.py
index 279195ca5..6fc0b8377 100644
--- a/numba_cuda/numba/cuda/core/entrypoints.py
+++ b/numba_cuda/numba/cuda/core/entrypoints.py
@@ -17,7 +17,7 @@ def init_all():
     If extensions have already been initialized, this function does nothing.
     """
     if _HAS_NUMBA:
-        from numba.core import entrypoints  # compat-ignore
+        from numba.core import entrypoints
 
         entrypoints.init_all()
 
diff --git a/numba_cuda/numba/cuda/core/inline_closurecall.py b/numba_cuda/numba/cuda/core/inline_closurecall.py
index 66a51b8d5..d1cd07f0e 100644
--- a/numba_cuda/numba/cuda/core/inline_closurecall.py
+++ b/numba_cuda/numba/cuda/core/inline_closurecall.py
@@ -5,6 +5,7 @@
 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.cuda.core import errors
@@ -235,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
@@ -1043,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/ir.py b/numba_cuda/numba/cuda/core/ir.py
index 2cdc74961..5951cd4b5 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,6 +14,10 @@
 from functools import total_ordering
 from io import StringIO
 
+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
diff --git a/numba_cuda/numba/cuda/core/ir_utils.py b/numba_cuda/numba/cuda/core/ir_utils.py
index ebe9f348a..b59f7fa48 100644
--- a/numba_cuda/numba/cuda/core/ir_utils.py
+++ b/numba_cuda/numba/cuda/core/ir_utils.py
@@ -10,6 +10,7 @@
 import warnings
 
 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
@@ -829,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
 
@@ -837,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):
diff --git a/numba_cuda/numba/cuda/core/sigutils.py b/numba_cuda/numba/cuda/core/sigutils.py
index 897275040..bd3925094 100644
--- a/numba_cuda/numba/cuda/core/sigutils.py
+++ b/numba_cuda/numba/cuda/core/sigutils.py
@@ -4,7 +4,7 @@
 from numba.cuda import types, typing, _HAS_NUMBA
 
 if _HAS_NUMBA:
-    from numba.core.typing import Signature as CoreSignature  # compat-ignore
+    from numba.core.typing import Signature as CoreSignature
 
 
 def is_signature(sig):
diff --git a/numba_cuda/numba/cuda/core/targetconfig.py b/numba_cuda/numba/cuda/core/targetconfig.py
index ac31fcd28..3473ec2ea 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/extending.py b/numba_cuda/numba/cuda/extending.py
index 9ce1f6ad5..11de8d096 100644
--- a/numba_cuda/numba/cuda/extending.py
+++ b/numba_cuda/numba/cuda/extending.py
@@ -219,12 +219,12 @@ def decorate(overload_func):
         # For generic/CPU targets, also register in numba.core registry
         if target in ("generic", "cpu"):
             try:
-                from numba.core.typing.templates import (  # compat-ignore
-                    make_overload_template as core_make_overload_template,  # compat-ignore
-                    infer as core_infer,  # compat-ignore
-                    infer_global as core_infer_global,  # compat-ignore
-                )  # compat-ignore
-                from numba.core import types as core_types  # compat-ignore
+                from numba.core.typing.templates import (
+                    make_overload_template as core_make_overload_template,
+                    infer as core_infer,
+                    infer_global as core_infer_global,
+                )
+                from numba.core import types as core_types
 
                 core_template = core_make_overload_template(
                     func,
@@ -310,10 +310,10 @@ def decorate(overload_func):
         # For generic/CPU targets, also register in numba.core registry
         if target in ("generic", "cpu"):
             try:
-                from numba.core.typing.templates import (  # compat-ignore
-                    make_overload_attribute_template as core_make_overload_attribute_template,  # compat-ignore
-                    infer_getattr as core_infer_getattr,  # compat-ignore
-                )  # compat-ignore
+                from numba.core.typing.templates import (
+                    make_overload_attribute_template as core_make_overload_attribute_template,
+                    infer_getattr as core_infer_getattr,
+                )
 
                 core_template = core_make_overload_attribute_template(
                     typ, attr, overload_func, **kwargs
@@ -348,10 +348,10 @@ def decorate(overload_func):
         target = kwargs.get("target", "cuda")
         if target in ("generic", "cpu"):
             try:
-                from numba.core.typing.templates import (  # compat-ignore
-                    make_overload_method_template as core_make_overload_method_template,  # compat-ignore
-                    infer_getattr as core_infer_getattr,  # compat-ignore
-                )  # compat-ignore
+                from numba.core.typing.templates import (
+                    make_overload_method_template as core_make_overload_method_template,
+                    infer_getattr as core_infer_getattr,
+                )
 
                 copied_kwargs2 = kwargs.copy()
                 core_template = core_make_overload_method_template(
diff --git a/numba_cuda/numba/cuda/lowering.py b/numba_cuda/numba/cuda/lowering.py
index 2bf11362a..8d787e408 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 (
@@ -1880,7 +1880,14 @@ def _lit_or_omitted(value):
     """Returns a Literal instance if the type of value is supported;
     otherwise, return `Omitted(value)`.
     """
+    excepts = LiteralTypingError
+    if _HAS_NUMBA:
+        from numba.core.errors import (
+            LiteralTypingError as CoreLiteralTypingError,
+        )
+
+        excepts = (LiteralTypingError, CoreLiteralTypingError)
     try:
         return types.literal(value)
-    except (LiteralTypingError, numba.core.errors.LiteralTypingError):
+    except excepts:
         return types.Omitted(value)
diff --git a/numba_cuda/numba/cuda/misc/mergesort.py b/numba_cuda/numba/cuda/misc/mergesort.py
index 79758485f..c6ed91993 100644
--- a/numba_cuda/numba/cuda/misc/mergesort.py
+++ b/numba_cuda/numba/cuda/misc/mergesort.py
@@ -23,30 +23,24 @@
 )
 
 
-def make_mergesort_impl(wrap, lt=None, is_argsort=False):
-    kwargs_lite = dict(no_cpython_wrapper=True, _nrt=False)
-
+def make_mergesort_impl(lt=None, is_argsort=False):
     # The less than
     if lt is None:
 
-        @wrap(**kwargs_lite)
         def lt(a, b):
             return a < b
     else:
-        lt = wrap(**kwargs_lite)(lt)
+        lt = lt
 
     if is_argsort:
 
-        @wrap(**kwargs_lite)
         def lessthan(a, b, vals):
             return lt(vals[a], vals[b])
     else:
 
-        @wrap(**kwargs_lite)
         def lessthan(a, b, vals):
             return lt(a, b)
 
-    @wrap(**kwargs_lite)
     def argmergesort_inner(arr, vals, ws):
         """The actual mergesort function
 
@@ -108,14 +102,12 @@ def argmergesort_inner(arr, vals, ws):
 
     # The top-level entry points
 
-    @wrap(no_cpython_wrapper=True)
     def mergesort(arr):
         "Inplace"
         ws = np.empty(arr.size // 2, dtype=arr.dtype)
         argmergesort_inner(arr, None, ws)
         return arr
 
-    @wrap(no_cpython_wrapper=True)
     def argmergesort(arr):
         "Out-of-place"
         idxs = np.arange(arr.size)
@@ -129,8 +121,4 @@ def argmergesort(arr):
 
 
 def make_jit_mergesort(*args, **kwargs):
-    from numba import njit
-
-    # NOTE: wrap with njit to allow recursion
-    #       because @register_jitable => @overload doesn't support recursion
-    return make_mergesort_impl(njit, *args, **kwargs)
+    return make_mergesort_impl(*args, **kwargs)
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 dcfc6bff4..7604301a1 100644
--- a/numba_cuda/numba/cuda/random.py
+++ b/numba_cuda/numba/cuda/random.py
@@ -3,7 +3,7 @@
 
 import math
 
-import numba.cuda as cuda
+from numba import cuda, jit
 from numba.cuda import (
     float32,
     float64,
@@ -13,7 +13,6 @@
     from_dtype,
 )
 from numba.cuda import config
-from numba import jit  # compat-ignore
 
 import numpy as np
 
diff --git a/numba_cuda/numba/cuda/target.py b/numba_cuda/numba/cuda/target.py
index ac3d875b5..bcfe09a23 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)
@@ -226,7 +222,7 @@ def load_additional_registries(self):
         # Install only implementations that are defined outside of numba (i.e.,
         # in third-party extensions) from Numba's builtin_registry.
         if importlib.util.find_spec("numba.core.imputils") is not None:
-            from numba.core.imputils import builtin_registry  # compat-ignore
+            from numba.core.imputils import builtin_registry
 
             self.install_external_registry(builtin_registry)
 
diff --git a/numba_cuda/numba/cuda/tests/core/serialize_usecases.py b/numba_cuda/numba/cuda/tests/core/serialize_usecases.py
index bd0fcfc51..a97db4af3 100644
--- a/numba_cuda/numba/cuda/tests/core/serialize_usecases.py
+++ b/numba_cuda/numba/cuda/tests/core/serialize_usecases.py
@@ -11,7 +11,10 @@
 import numpy as np
 import numpy.random as nprand
 
-from numba import jit
+from numba.cuda import _HAS_NUMBA
+
+if _HAS_NUMBA:
+    from numba import jit
 
 
 @jit("int32(int32, int32)")
diff --git a/numba_cuda/numba/cuda/tests/core/test_serialize.py b/numba_cuda/numba/cuda/tests/core/test_serialize.py
index b5a49d2d6..94683fe56 100644
--- a/numba_cuda/numba/cuda/tests/core/test_serialize.py
+++ b/numba_cuda/numba/cuda/tests/core/test_serialize.py
@@ -11,26 +11,19 @@
 from multiprocessing import get_context
 
 from numba.cuda import _HAS_NUMBA
+import numba
 
 if _HAS_NUMBA:
-    from numba.core.errors import TypingError  # compat-ignore
-    import numba  # compat-ignore
+    from numba.core.errors import TypingError
+    from numba.core.target_extension import resolve_dispatcher_from_str
 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
 
-try:
-    from numba.core.target_extension import resolve_dispatcher_from_str
-except ImportError:
-    resolve_dispatcher_from_str = None
-
 
-@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):
@@ -66,6 +59,7 @@ def check_result(func):
             new_func = pickle.loads(pickled)
             check_result(new_func)
 
+    @skip_on_standalone_numba_cuda
     def test_call_with_sig(self):
         from .serialize_usecases import add_with_sig
 
@@ -73,6 +67,7 @@ def test_call_with_sig(self):
         # Compilation has been disabled => float inputs will be coerced to int
         self.run_with_protocols(self.check_call, add_with_sig, 5, (1.2, 4.2))
 
+    @skip_on_standalone_numba_cuda
     def test_call_without_sig(self):
         from .serialize_usecases import add_without_sig
 
@@ -85,6 +80,7 @@ def test_call_without_sig(self):
             self.check_call, add_without_sig, "abc", ("a", "bc")
         )
 
+    @skip_on_standalone_numba_cuda
     def test_call_nopython(self):
         from .serialize_usecases import add_nopython
 
@@ -94,6 +90,7 @@ def test_call_nopython(self):
             self.check_call, add_nopython, TypingError, (object(), object())
         )
 
+    @skip_on_standalone_numba_cuda
     def test_call_nopython_fail(self):
         from .serialize_usecases import add_nopython_fail
 
@@ -102,6 +99,7 @@ def test_call_nopython_fail(self):
             self.check_call, add_nopython_fail, TypingError, (1, 2)
         )
 
+    @skip_on_standalone_numba_cuda
     def test_call_objmode_with_global(self):
         from .serialize_usecases import get_global_objmode
 
@@ -109,12 +107,14 @@ def test_call_objmode_with_global(self):
             self.check_call, get_global_objmode, 7.5, (2.5,)
         )
 
+    @skip_on_standalone_numba_cuda
     def test_call_closure(self):
         from .serialize_usecases import closure
 
         inner = closure(1)
         self.run_with_protocols(self.check_call, inner, 6, (2, 3))
 
+    @skip_on_standalone_numba_cuda
     def check_call_closure_with_globals(self, **jit_args):
         from .serialize_usecases import closure_with_globals
 
@@ -127,30 +127,35 @@ def test_call_closure_with_globals_nopython(self):
     def test_call_closure_with_globals_objmode(self):
         self.check_call_closure_with_globals(forceobj=True)
 
+    @skip_on_standalone_numba_cuda
     def test_call_closure_calling_other_function(self):
         from .serialize_usecases import closure_calling_other_function
 
         inner = closure_calling_other_function(3.0)
         self.run_with_protocols(self.check_call, inner, 11.0, (4.0, 6.0))
 
+    @skip_on_standalone_numba_cuda
     def test_call_closure_calling_other_closure(self):
         from .serialize_usecases import closure_calling_other_closure
 
         inner = closure_calling_other_closure(3.0)
         self.run_with_protocols(self.check_call, inner, 8.0, (4.0,))
 
+    @skip_on_standalone_numba_cuda
     def test_call_dyn_func(self):
         from .serialize_usecases import dyn_func
 
         # Check serializing a dynamically-created function
         self.run_with_protocols(self.check_call, dyn_func, 36, (6,))
 
+    @skip_on_standalone_numba_cuda
     def test_call_dyn_func_objmode(self):
         from .serialize_usecases import dyn_func_objmode
 
         # Same with an object mode function
         self.run_with_protocols(self.check_call, dyn_func_objmode, 36, (6,))
 
+    @skip_on_standalone_numba_cuda
     def test_renamed_module(self):
         from .serialize_usecases import get_renamed_module
 
@@ -161,6 +166,7 @@ def test_renamed_module(self):
             self.check_call, get_renamed_module, expected, (0.0,)
         )
 
+    @skip_on_standalone_numba_cuda
     def test_other_process(self):
         """
         Check that reconstructing doesn't depend on resources already
@@ -180,6 +186,7 @@ def test_other_process(self):
             """.format(**locals())
         subprocess.check_call([sys.executable, "-c", code])
 
+    @skip_on_standalone_numba_cuda
     def test_reuse(self):
         """
         Check that deserializing the same function multiple times re-uses
@@ -219,6 +226,7 @@ def test_reuse(self):
         g.disable_compile()
         self.assertEqual(g(2, 4), 13)
 
+    @skip_on_standalone_numba_cuda
     def test_imp_deprecation(self):
         """
         The imp module was deprecated in v3.4 in favour of importlib
@@ -312,7 +320,7 @@ class Klass:
 
     @skip_on_standalone_numba_cuda
     def test_dynamic_class_issue_7356(self):
-        import numba  # compat-ignore
+        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..adba8252b 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 e90ed929d..b21011dbf 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_array_alignment.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_array_alignment.py
@@ -4,11 +4,11 @@
 import re
 import itertools
 import numpy as np
-import numba.cuda as cuda
+from numba import cuda
 from numba.cuda import _HAS_NUMBA
 
 if _HAS_NUMBA:
-    from numba.core.errors import TypingError  # compat-ignore
+    from numba.core.errors import TypingError
 else:
     from numba.cuda.core.errors import TypingError
 from numba.cuda.testing import (
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_datetime.py b/numba_cuda/numba/cuda/tests/cudapy/test_datetime.py
index e44522416..07ad931d8 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_datetime.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_datetime.py
@@ -4,7 +4,7 @@
 import numpy as np
 
 from numba.cuda import vectorize, guvectorize
-import numba.cuda as cuda
+from numba import cuda
 from numba.cuda.np.numpy_support import from_dtype
 from numba.cuda.testing import CUDATestCase, skip_on_cudasim
 import unittest
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..b2ba8d9a8 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 e1e035c3c..9505ae912 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_enums.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_enums.py
@@ -12,8 +12,8 @@
 from numba.cuda import _HAS_NUMBA
 
 if _HAS_NUMBA:
-    from numba import njit  # compat-ignore
-import numba.cuda as cuda
+    from numba import njit
+from numba import cuda
 from numba.cuda import types
 from numba.cuda.testing import (
     unittest,
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_extending.py b/numba_cuda/numba/cuda/tests/cudapy/test_extending.py
index f0e920b83..4c1b56b85 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 019482321..223b9ad2c 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_gufunc.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_gufunc.py
@@ -6,7 +6,7 @@
 from collections import namedtuple
 from numba.cuda import void, int32, float32, float64
 from numba.cuda import guvectorize
-import numba.cuda as cuda
+from numba import cuda
 from numba.cuda.testing import skip_on_cudasim, CUDATestCase
 import unittest
 import warnings
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 7cb4ddca3..9a16a631c 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_gufunc_scalar.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_gufunc_scalar.py
@@ -9,7 +9,7 @@
 
 import numpy as np
 from numba.cuda import guvectorize
-import numba.cuda as cuda
+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_intrinsics.py b/numba_cuda/numba/cuda/tests/cudapy/test_intrinsics.py
index d8490cefe..9072ff50f 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_intrinsics.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_intrinsics.py
@@ -10,7 +10,7 @@
 from numba.cuda import _HAS_NUMBA
 
 if _HAS_NUMBA:
-    from numba.core.errors import TypingError  # compat-ignore
+    from numba.core.errors import TypingError
 else:
     from numba.cuda.core.errors import TypingError
 from numba.cuda.types import f2
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_ipc.py b/numba_cuda/numba/cuda/tests/cudapy/test_ipc.py
index 4f27b3a88..3f21dd7bd 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_ipc.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_ipc.py
@@ -10,7 +10,7 @@
 
 import numpy as np
 
-import numba.cuda as cuda
+from numba import cuda
 from numba.cuda.testing import (
     skip_on_arm,
     skip_on_cudasim,
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_ir.py b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
index f7ff4195f..994a9168c 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
@@ -9,8 +9,8 @@
 from numba.cuda import _HAS_NUMBA
 
 if _HAS_NUMBA:
-    from numba import objmode  # compat-ignore
-    from numba import njit  # compat-ignore
+    from numba import objmode
+    from numba import njit
 from numba.cuda.core import ir
 from numba.cuda import compiler
 from numba.cuda.core import errors
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_localmem.py b/numba_cuda/numba/cuda/tests/cudapy/test_localmem.py
index 4465223e3..25cc4b1ef 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_localmem.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_localmem.py
@@ -9,7 +9,7 @@
 from numba.cuda import _HAS_NUMBA
 
 if _HAS_NUMBA:
-    from numba.core.errors import TypingError  # compat-ignore
+    from numba.core.errors import TypingError
 else:
     from numba.cuda.core.errors import TypingError
 from numba.cuda.testing import unittest, CUDATestCase, skip_on_cudasim
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_overload.py b/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
index ff2e3a89b..d1fddbd81 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
@@ -1,19 +1,19 @@
 # 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
 import numba.cuda.types as types
 from numba.cuda import _HAS_NUMBA
 
 if _HAS_NUMBA:
-    from numba.core.errors import TypingError  # compat-ignore
-    from numba import njit  # compat-ignore
-    import numba  # compat-ignore
+    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  # compat-ignore
+from numba.core.typing.typeof import typeof as cpu_typeof
 from numba.cuda.testing import (
     CUDATestCase,
     skip_on_cudasim,
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_serialize.py b/numba_cuda/numba/cuda/tests/cudapy/test_serialize.py
index 0429b0013..8c126013d 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_serialize.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_serialize.py
@@ -4,7 +4,7 @@
 import pickle
 import numpy as np
 from numba.cuda import vectorize
-import numba.cuda as cuda
+from numba import cuda
 from numba.cuda import types
 from numba.cuda.testing import skip_on_cudasim, CUDATestCase
 import unittest
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_sm.py b/numba_cuda/numba/cuda/tests/cudapy/test_sm.py
index a58922da6..75fa8a941 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_sm.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_sm.py
@@ -6,7 +6,7 @@
 from numba.cuda import _HAS_NUMBA
 
 if _HAS_NUMBA:
-    from numba.core.errors import TypingError  # compat-ignore
+    from numba.core.errors import TypingError
     from numba.cuda.core.errors import TypingError as cudaTypingError
 else:
     from numba.cuda.core.errors import TypingError as cudaTypingError
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 6b3d156d5..0fc044b7f 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_sm_creation.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_sm_creation.py
@@ -7,7 +7,7 @@
 from numba.cuda import _HAS_NUMBA
 
 if _HAS_NUMBA:
-    from numba.core.errors import TypingError  # compat-ignore
+    from numba.core.errors import TypingError
 else:
     from numba.cuda.core.errors import TypingError
 from numba.cuda.testing import unittest, CUDATestCase
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_ufuncs.py b/numba_cuda/numba/cuda/tests/cudapy/test_ufuncs.py
index f09205c1f..c36101706 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/cudapy/test_vectorize.py b/numba_cuda/numba/cuda/tests/cudapy/test_vectorize.py
index 9b779e248..3a41b1234 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_vectorize.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_vectorize.py
@@ -7,7 +7,7 @@
 from functools import partial
 from itertools import product
 from numba.cuda import vectorize as cuda_vectorize
-from numba import cuda, vectorize as numba_vectorize  # compat-ignore
+from numba import cuda, vectorize as numba_vectorize
 from numba.cuda.types import int32, float32, float64
 from numba.cuda.cudadrv.driver import CudaAPIError, driver
 from numba.cuda.testing import skip_on_cudasim
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_warning.py b/numba_cuda/numba/cuda/tests/cudapy/test_warning.py
index 3d2b1eb3e..a80278eaf 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_warning.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_warning.py
@@ -2,7 +2,7 @@
 # SPDX-License-Identifier: BSD-2-Clause
 
 import numpy as np
-import numba.cuda as cuda
+from numba import cuda
 from numba.cuda.cudadrv import driver
 from numba.cuda.testing import (
     unittest,
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/nrt/test_nrt.py b/numba_cuda/numba/cuda/tests/nrt/test_nrt.py
index b4a45d7a1..1efeed919 100644
--- a/numba_cuda/numba/cuda/tests/nrt/test_nrt.py
+++ b/numba_cuda/numba/cuda/tests/nrt/test_nrt.py
@@ -12,7 +12,7 @@
 from numba.cuda.cudadrv.nvrtc import compile
 from numba.cuda import types
 from numba.cuda.typing import signature
-import numba.cuda as cuda
+from numba import cuda
 from numba.cuda import config
 from numba.cuda.typing.templates import AbstractTemplate
 from numba.cuda.cudadrv.linkable_code import (
@@ -218,7 +218,7 @@ def tearDown(self):
     def test_stats_env_var_explicit_on(self):
         # Checks that explicitly turning the stats on via the env var works.
         src = """if 1:
-        import numba.cuda as cuda
+        from numba import cuda
         from numba.cuda.memory_management import rtsys
         import numpy as np
 
@@ -256,7 +256,7 @@ def foo():
 
     def check_env_var_off(self, env):
         src = """if 1:
-        import numba.cuda as cuda
+        from numba import cuda
         import numpy as np
         from numba.cuda.memory_management import rtsys
 
diff --git a/numba_cuda/numba/cuda/tests/support.py b/numba_cuda/numba/cuda/tests/support.py
index b9ad46c18..31d5ba41b 100644
--- a/numba_cuda/numba/cuda/tests/support.py
+++ b/numba_cuda/numba/cuda/tests/support.py
@@ -41,8 +41,8 @@
 if _HAS_NUMBA:
     from numba.core.extending import (
         typeof_impl as upstream_typeof_impl,
-    )  # compat-ignore
-    from numba.core import types as upstream_types  # compat-ignore
+    )
+    from numba.core import types as upstream_types
 else:
     upstream_typeof_impl = None
     upstream_types = None
diff --git a/numba_cuda/numba/cuda/tests/test_array_reductions.py b/numba_cuda/numba/cuda/tests/test_array_reductions.py
index 06ab07a30..2823a00cd 100644
--- a/numba_cuda/numba/cuda/tests/test_array_reductions.py
+++ b/numba_cuda/numba/cuda/tests/test_array_reductions.py
@@ -2,7 +2,7 @@
 # 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.misc.special import literal_unroll
 from numba.cuda import config
diff --git a/numba_cuda/numba/cuda/tests/test_extending_types.py b/numba_cuda/numba/cuda/tests/test_extending_types.py
index 0c78efddf..02d752602 100644
--- a/numba_cuda/numba/cuda/tests/test_extending_types.py
+++ b/numba_cuda/numba/cuda/tests/test_extending_types.py
@@ -10,7 +10,7 @@
 from numba.cuda import _HAS_NUMBA
 
 if _HAS_NUMBA:
-    from numba.core.errors import TypingError, NumbaTypeError  # compat-ignore
+    from numba.core.errors import TypingError, NumbaTypeError
 else:
     from numba.cuda.core.errors import TypingError
 from numba.cuda.extending import make_attribute_wrapper
diff --git a/numba_cuda/numba/cuda/typing/templates.py b/numba_cuda/numba/cuda/typing/templates.py
index 3f78298c6..e9078f6e5 100644
--- a/numba_cuda/numba/cuda/typing/templates.py
+++ b/numba_cuda/numba/cuda/typing/templates.py
@@ -28,7 +28,7 @@
 from numba.cuda import _HAS_NUMBA
 
 if _HAS_NUMBA:
-    from numba.core.typing import Signature as CoreSignature  # compat-ignore
+    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")
diff --git a/numba_cuda/numba/cuda/typing/typeof.py b/numba_cuda/numba/cuda/typing/typeof.py
index 3c718c9d4..caece355c 100644
--- a/numba_cuda/numba/cuda/typing/typeof.py
+++ b/numba_cuda/numba/cuda/typing/typeof.py
@@ -8,7 +8,7 @@
 
 import numpy as np
 from numpy.random.bit_generator import BitGenerator
-
+from numba.cuda import _HAS_NUMBA
 from numba.cuda import types
 from numba.cuda.core import errors
 from numba.cuda import utils
@@ -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):

From caaa5ccb2d0ce947ce0cadab0d8b30f806273c04 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Thu, 30 Oct 2025 10:34:13 -0700
Subject: [PATCH 16/25] cleaning up some of the guards, removing some of the
 standalone_numba_cuda reasons (not needed), cleaning up some imports

---
 numba_cuda/numba/cuda/__init__.py                | 11 +----------
 numba_cuda/numba/cuda/core/errors.py             |  4 ++--
 numba_cuda/numba/cuda/lowering.py                |  6 +++---
 .../numba/cuda/tests/core/serialize_usecases.py  |  6 ++----
 .../numba/cuda/tests/cudapy/test_datetime.py     |  3 +--
 numba_cuda/numba/cuda/tests/cudapy/test_enums.py |  7 ++-----
 .../cuda/tests/cudapy/test_gufunc_scalar.py      |  3 +--
 numba_cuda/numba/cuda/tests/cudapy/test_math.py  |  3 +--
 .../numba/cuda/tests/cudapy/test_overload.py     |  8 ++------
 .../numba/cuda/tests/cudapy/test_serialize.py    |  3 +--
 numba_cuda/numba/cuda/tests/cudapy/test_sm.py    | 16 +++++++++-------
 numba_cuda/numba/cuda/tests/support.py           |  5 +----
 .../numba/cuda/tests/test_extending_types.py     |  2 +-
 13 files changed, 27 insertions(+), 50 deletions(-)

diff --git a/numba_cuda/numba/cuda/__init__.py b/numba_cuda/numba/cuda/__init__.py
index f7ed99364..7ba0e5594 100644
--- a/numba_cuda/numba/cuda/__init__.py
+++ b/numba_cuda/numba/cuda/__init__.py
@@ -7,10 +7,8 @@
 import warnings
 import sys
 
-# Import version from the parent package
 from numba_cuda._version import __version__
 
-# Generate version_info early to avoid circular import issues
 from numba_cuda._version import generate_version_info
 
 version_info = generate_version_info(__version__)
@@ -22,14 +20,7 @@
 # Re-export all type names
 from numba.cuda.types import *
 
-# Check if upstream numba is available - needs to be defined early to avoid circular imports
-try:
-    import numba
-
-    _HAS_NUMBA = True
-except ImportError:
-    _HAS_NUMBA = False
-
+_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/core/errors.py b/numba_cuda/numba/cuda/core/errors.py
index 68c63f851..0d15e3c6e 100644
--- a/numba_cuda/numba/cuda/core/errors.py
+++ b/numba_cuda/numba/cuda/core/errors.py
@@ -15,7 +15,7 @@
 from functools import wraps
 from abc import abstractmethod
 
-from numba_cuda._version import __version__
+import numba_cuda
 
 # Filled at the end
 __all__ = []
@@ -456,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
-""" % (__version__, feedback_details)
+""" % (numba_cuda.__version__, feedback_details)
 
 error_extras = dict()
 error_extras["unsupported_error"] = unsupported_error_info
diff --git a/numba_cuda/numba/cuda/lowering.py b/numba_cuda/numba/cuda/lowering.py
index 8d787e408..70b3e8524 100644
--- a/numba_cuda/numba/cuda/lowering.py
+++ b/numba_cuda/numba/cuda/lowering.py
@@ -1880,14 +1880,14 @@ def _lit_or_omitted(value):
     """Returns a Literal instance if the type of value is supported;
     otherwise, return `Omitted(value)`.
     """
-    excepts = LiteralTypingError
+    typing_errors = LiteralTypingError
     if _HAS_NUMBA:
         from numba.core.errors import (
             LiteralTypingError as CoreLiteralTypingError,
         )
 
-        excepts = (LiteralTypingError, CoreLiteralTypingError)
+        typing_errors = (LiteralTypingError, CoreLiteralTypingError)
     try:
         return types.literal(value)
-    except excepts:
+    except typing_errors:
         return types.Omitted(value)
diff --git a/numba_cuda/numba/cuda/tests/core/serialize_usecases.py b/numba_cuda/numba/cuda/tests/core/serialize_usecases.py
index a97db4af3..37e676b6e 100644
--- a/numba_cuda/numba/cuda/tests/core/serialize_usecases.py
+++ b/numba_cuda/numba/cuda/tests/core/serialize_usecases.py
@@ -11,10 +11,8 @@
 import numpy as np
 import numpy.random as nprand
 
-from numba.cuda import _HAS_NUMBA
-
-if _HAS_NUMBA:
-    from numba import jit
+# This does not need a guard, it's already guarded at the import site
+from numba import jit
 
 
 @jit("int32(int32, int32)")
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_datetime.py b/numba_cuda/numba/cuda/tests/cudapy/test_datetime.py
index 07ad931d8..2c513dd0a 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_datetime.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_datetime.py
@@ -3,8 +3,7 @@
 
 import numpy as np
 
-from numba.cuda import vectorize, guvectorize
-from numba import cuda
+from numba import cuda, vectorize, guvectorize
 from numba.cuda.np.numpy_support import from_dtype
 from numba.cuda.testing import CUDATestCase, skip_on_cudasim
 import unittest
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_enums.py b/numba_cuda/numba/cuda/tests/cudapy/test_enums.py
index 9505ae912..c5c5d4a2a 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_enums.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_enums.py
@@ -8,8 +8,7 @@
 import numpy as np
 
 from numba.cuda import int16, int32
-from numba.cuda import vectorize
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import vectorize, _HAS_NUMBA
 
 if _HAS_NUMBA:
     from numba import njit
@@ -67,9 +66,7 @@ def f(out):
         f(expected)
         self.assertPreciseEqual(expected, got)
 
-    @skip_on_standalone_numba_cuda(
-        "Test not supported in standalone numba_cuda"
-    )
+    @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_gufunc_scalar.py b/numba_cuda/numba/cuda/tests/cudapy/test_gufunc_scalar.py
index 9a16a631c..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,8 +8,7 @@
 """
 
 import numpy as np
-from numba.cuda import guvectorize
-from numba import 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_math.py b/numba_cuda/numba/cuda/tests/cudapy/test_math.py
index 97c2a9613..71b089cf4 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_math.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_math.py
@@ -10,8 +10,7 @@
     skip_on_cudasim,
 )
 from numba.cuda.np import numpy_support
-from numba import vectorize
-import numba.cuda as cuda
+from numba import cuda, vectorize
 from numba.cuda import float32, float64, int32, void, int64
 import math
 
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_overload.py b/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
index d1fddbd81..6a124a063 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
@@ -237,9 +237,7 @@ def check_overload(self, kernel, expected):
         cuda.jit(kernel)[1, 1](x)
         self.assertEqual(x[0], expected)
 
-    @skip_on_standalone_numba_cuda(
-        "Test not supported in standalone numba_cuda"
-    )
+    @skip_on_standalone_numba_cuda
     def check_overload_cpu(self, kernel, expected):
         x = np.ones(1, dtype=np.int32)
         njit(kernel)(x)
@@ -344,9 +342,7 @@ def kernel(x):
         expected = GENERIC_TARGET_OL_CALLS_TARGET_OL * GENERIC_TARGET_OL
         self.check_overload_cpu(kernel, expected)
 
-    @skip_on_standalone_numba_cuda(
-        "Test not supported in standalone numba_cuda"
-    )
+    @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)
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_serialize.py b/numba_cuda/numba/cuda/tests/cudapy/test_serialize.py
index 8c126013d..f82abfc67 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_serialize.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_serialize.py
@@ -3,8 +3,7 @@
 
 import pickle
 import numpy as np
-from numba.cuda import vectorize
-from numba import cuda
+from numba import cuda, vectorize
 from numba.cuda import types
 from numba.cuda.testing import skip_on_cudasim, CUDATestCase
 import unittest
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_sm.py b/numba_cuda/numba/cuda/tests/cudapy/test_sm.py
index 75fa8a941..36d1e64db 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_sm.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_sm.py
@@ -6,10 +6,8 @@
 from numba.cuda import _HAS_NUMBA
 
 if _HAS_NUMBA:
-    from numba.core.errors import TypingError
-    from numba.cuda.core.errors import TypingError as cudaTypingError
-else:
-    from numba.cuda.core.errors import TypingError as cudaTypingError
+    from numba.core.errors import TypingError as NumbaTypingError
+from numba.cuda.core.errors import TypingError
 from numba.cuda import types
 from numba.cuda.testing import unittest, CUDATestCase, skip_on_cudasim
 
@@ -407,7 +405,7 @@ def test_invalid_array_type(self):
         def unsupported_type():
             arr = cuda.shared.array(10, dtype=np.dtype("O"))  # noqa: F841
 
-        with self.assertRaisesRegex(cudaTypingError, rgx):
+        with self.assertRaisesRegex(TypingError, rgx):
             cuda.jit(void())(unsupported_type)
 
         rgx = ".*Invalid NumPy dtype specified: 'int33'.*"
@@ -415,8 +413,12 @@ def unsupported_type():
         def invalid_string_type():
             arr = cuda.shared.array(10, dtype="int33")  # noqa: F841
 
-        with self.assertRaisesRegex(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/support.py b/numba_cuda/numba/cuda/tests/support.py
index 31d5ba41b..e60fb969d 100644
--- a/numba_cuda/numba/cuda/tests/support.py
+++ b/numba_cuda/numba/cuda/tests/support.py
@@ -43,9 +43,6 @@
         typeof_impl as upstream_typeof_impl,
     )
     from numba.core import types as upstream_types
-else:
-    upstream_typeof_impl = None
-    upstream_types = None
 
 
 class EnableNRTStatsMixin(object):
@@ -773,7 +770,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_extending_types.py b/numba_cuda/numba/cuda/tests/test_extending_types.py
index 02d752602..cee9222a3 100644
--- a/numba_cuda/numba/cuda/tests/test_extending_types.py
+++ b/numba_cuda/numba/cuda/tests/test_extending_types.py
@@ -12,7 +12,7 @@
 if _HAS_NUMBA:
     from numba.core.errors import TypingError, NumbaTypeError
 else:
-    from numba.cuda.core.errors import TypingError
+    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

From 31af2441542020ae7f89be302b66166d42c497bd Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Thu, 30 Oct 2025 11:05:56 -0700
Subject: [PATCH 17/25] skip array reductions test on sim

---
 numba_cuda/numba/cuda/tests/test_array_reductions.py | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/numba_cuda/numba/cuda/tests/test_array_reductions.py b/numba_cuda/numba/cuda/tests/test_array_reductions.py
index 2823a00cd..167095e1a 100644
--- a/numba_cuda/numba/cuda/tests/test_array_reductions.py
+++ b/numba_cuda/numba/cuda/tests/test_array_reductions.py
@@ -4,10 +4,12 @@
 
 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.

From 15dc22a709d8f56eb581219aa788e3565b521abe Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Thu, 30 Oct 2025 11:10:55 -0700
Subject: [PATCH 18/25] add an extra guard in test_overload for cpu tests

---
 numba_cuda/numba/cuda/tests/cudapy/test_overload.py | 9 +++++++--
 1 file changed, 7 insertions(+), 2 deletions(-)

diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_overload.py b/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
index 6a124a063..9fbab1590 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
@@ -2,7 +2,7 @@
 # SPDX-License-Identifier: BSD-2-Clause
 
 from numba import cuda
-import numba.cuda.types as types
+from numba.cuda import types
 from numba.cuda import _HAS_NUMBA
 
 if _HAS_NUMBA:
@@ -338,7 +338,12 @@ 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)
 

From be785c0b6a51d4067545b9def46fc7e7cbc541d4 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Thu, 30 Oct 2025 12:30:28 -0700
Subject: [PATCH 19/25] rename _HAS_NUMBA to HAS_NUMBA, add guard around cpu
 jiting in random

---
 numba_cuda/numba/cuda/__init__.py                      |  2 +-
 numba_cuda/numba/cuda/core/entrypoints.py              |  4 ++--
 numba_cuda/numba/cuda/core/inline_closurecall.py       |  4 ++--
 numba_cuda/numba/cuda/core/ir.py                       | 10 +++++-----
 numba_cuda/numba/cuda/core/ir_utils.py                 |  4 ++--
 numba_cuda/numba/cuda/core/sigutils.py                 |  8 ++++----
 numba_cuda/numba/cuda/core/targetconfig.py             |  4 ++--
 numba_cuda/numba/cuda/lowering.py                      |  4 ++--
 numba_cuda/numba/cuda/random.py                        |  6 +++++-
 numba_cuda/numba/cuda/target.py                        |  4 ++--
 numba_cuda/numba/cuda/testing.py                       |  4 ++--
 numba_cuda/numba/cuda/tests/core/test_serialize.py     |  4 ++--
 .../numba/cuda/tests/cudapy/cache_with_cpu_usecases.py |  4 ++--
 .../numba/cuda/tests/cudapy/test_array_alignment.py    |  4 ++--
 numba_cuda/numba/cuda/tests/cudapy/test_device_func.py |  4 ++--
 numba_cuda/numba/cuda/tests/cudapy/test_enums.py       |  4 ++--
 numba_cuda/numba/cuda/tests/cudapy/test_extending.py   |  4 ++--
 numba_cuda/numba/cuda/tests/cudapy/test_intrinsics.py  |  4 ++--
 numba_cuda/numba/cuda/tests/cudapy/test_ir.py          |  8 ++++----
 numba_cuda/numba/cuda/tests/cudapy/test_localmem.py    |  4 ++--
 numba_cuda/numba/cuda/tests/cudapy/test_overload.py    |  6 +++---
 numba_cuda/numba/cuda/tests/cudapy/test_random.py      |  7 ++++++-
 numba_cuda/numba/cuda/tests/cudapy/test_sm.py          |  6 +++---
 numba_cuda/numba/cuda/tests/cudapy/test_sm_creation.py |  4 ++--
 numba_cuda/numba/cuda/tests/cudapy/test_ufuncs.py      |  4 ++--
 numba_cuda/numba/cuda/tests/support.py                 |  6 +++---
 numba_cuda/numba/cuda/tests/test_extending.py          |  4 ++--
 numba_cuda/numba/cuda/tests/test_extending_types.py    |  4 ++--
 numba_cuda/numba/cuda/typing/templates.py              |  8 ++++----
 numba_cuda/numba/cuda/typing/typeof.py                 |  4 ++--
 30 files changed, 78 insertions(+), 69 deletions(-)

diff --git a/numba_cuda/numba/cuda/__init__.py b/numba_cuda/numba/cuda/__init__.py
index 7ba0e5594..29ac69d68 100644
--- a/numba_cuda/numba/cuda/__init__.py
+++ b/numba_cuda/numba/cuda/__init__.py
@@ -20,7 +20,7 @@
 # Re-export all type names
 from numba.cuda.types import *
 
-_HAS_NUMBA = importlib.util.find_spec("numba") is not None
+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/core/entrypoints.py b/numba_cuda/numba/cuda/core/entrypoints.py
index 6fc0b8377..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
+from numba.cuda import HAS_NUMBA
 
 _already_initialized = False
 logger = logging.getLogger(__name__)
@@ -16,7 +16,7 @@ def init_all():
 
     If extensions have already been initialized, this function does nothing.
     """
-    if _HAS_NUMBA:
+    if HAS_NUMBA:
         from numba.core import entrypoints
 
         entrypoints.init_all()
diff --git a/numba_cuda/numba/cuda/core/inline_closurecall.py b/numba_cuda/numba/cuda/core/inline_closurecall.py
index d1cd07f0e..4b5dd09f7 100644
--- a/numba_cuda/numba/cuda/core/inline_closurecall.py
+++ b/numba_cuda/numba/cuda/core/inline_closurecall.py
@@ -5,7 +5,7 @@
 import copy
 import ctypes
 import numba.cuda.core.analysis
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 from numba.cuda import types, config, cgutils
 from numba.cuda.core import ir
 from numba.cuda.core import errors
@@ -236,7 +236,7 @@ def check_reduce_func(func_ir, func_var):
                             analysis"
         )
     if isinstance(reduce_func, (ir.FreeVar, ir.Global)):
-        if _HAS_NUMBA:
+        if HAS_NUMBA:
             from numba.core.registry import CPUDispatcher
 
             if not isinstance(reduce_func.value, CPUDispatcher):
diff --git a/numba_cuda/numba/cuda/core/ir.py b/numba_cuda/numba/cuda/core/ir.py
index 5951cd4b5..b907b2f14 100644
--- a/numba_cuda/numba/cuda/core/ir.py
+++ b/numba_cuda/numba/cuda/core/ir.py
@@ -14,9 +14,9 @@
 from functools import total_ordering
 from io import StringIO
 
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     import numba
 from numba.cuda.core import errors
 from numba.cuda.core import config
@@ -28,7 +28,7 @@
     ConstantInferenceError,
 )
 from numba.cuda.core import consts
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
 # terminal color markup
 _termcolor = errors.termcolor()
@@ -780,7 +780,7 @@ def __repr__(self):
 class Del(Stmt):
     def __init__(self, value, loc):
         assert isinstance(value, str)
-        if _HAS_NUMBA:
+        if HAS_NUMBA:
             assert isinstance(loc, (Loc, numba.core.ir.Loc))
         else:
             assert isinstance(loc)
@@ -1349,7 +1349,7 @@ class Block(EqualityCheckMixin):
     """A code block"""
 
     def __init__(self, scope, loc):
-        if _HAS_NUMBA:
+        if HAS_NUMBA:
             assert isinstance(scope, (Scope, numba.core.ir.Scope))
             assert isinstance(loc, (Loc, numba.core.ir.Loc))
         else:
diff --git a/numba_cuda/numba/cuda/core/ir_utils.py b/numba_cuda/numba/cuda/core/ir_utils.py
index b59f7fa48..225346985 100644
--- a/numba_cuda/numba/cuda/core/ir_utils.py
+++ b/numba_cuda/numba/cuda/core/ir_utils.py
@@ -10,7 +10,7 @@
 import warnings
 
 import numba.cuda
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 from numba.cuda import types
 from numba.cuda.core import ir
 from numba.cuda import typing
@@ -830,7 +830,7 @@ def has_no_side_effect(rhs, lives, call_table):
         ):
             return True
 
-        if _HAS_NUMBA:
+        if HAS_NUMBA:
             from numba.core.registry import CPUDispatcher
             from numba.cuda.np.linalg import dot_3_mv_check_args
 
diff --git a/numba_cuda/numba/cuda/core/sigutils.py b/numba_cuda/numba/cuda/core/sigutils.py
index bd3925094..cd7f9a65e 100644
--- a/numba_cuda/numba/cuda/core/sigutils.py
+++ b/numba_cuda/numba/cuda/core/sigutils.py
@@ -1,9 +1,9 @@
 # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 # SPDX-License-Identifier: BSD-2-Clause
 
-from numba.cuda import types, typing, _HAS_NUMBA
+from numba.cuda import types, typing, HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba.core.typing import Signature as CoreSignature
 
 
@@ -13,7 +13,7 @@ def is_signature(sig):
     specification (for user-facing APIs).
     """
     sig_types = (str, tuple, typing.Signature)
-    if _HAS_NUMBA:
+    if HAS_NUMBA:
         sig_types = (str, tuple, typing.Signature, CoreSignature)
     return isinstance(sig, sig_types)
 
@@ -42,7 +42,7 @@ def normalize_signature(sig):
         args, return_type = parsed, None
     else:
         sig_types = (typing.Signature,)
-        if _HAS_NUMBA:
+        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/targetconfig.py b/numba_cuda/numba/cuda/core/targetconfig.py
index 3473ec2ea..0559a130b 100644
--- a/numba_cuda/numba/cuda/core/targetconfig.py
+++ b/numba_cuda/numba/cuda/core/targetconfig.py
@@ -12,7 +12,7 @@
 
 from types import MappingProxyType
 from numba.cuda import utils
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
 
 class Option:
@@ -49,7 +49,7 @@ def doc(self):
         return self._doc
 
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba.core.targetconfig import ConfigStack, _FlagsStack
 else:
 
diff --git a/numba_cuda/numba/cuda/lowering.py b/numba_cuda/numba/cuda/lowering.py
index 70b3e8524..9df637fd4 100644
--- a/numba_cuda/numba/cuda/lowering.py
+++ b/numba_cuda/numba/cuda/lowering.py
@@ -8,7 +8,7 @@
 
 from llvmlite import ir as llvm_ir
 
-from numba.cuda import _HAS_NUMBA
+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 (
@@ -1881,7 +1881,7 @@ def _lit_or_omitted(value):
     otherwise, return `Omitted(value)`.
     """
     typing_errors = LiteralTypingError
-    if _HAS_NUMBA:
+    if HAS_NUMBA:
         from numba.core.errors import (
             LiteralTypingError as CoreLiteralTypingError,
         )
diff --git a/numba_cuda/numba/cuda/random.py b/numba_cuda/numba/cuda/random.py
index 7604301a1..8cf7cc079 100644
--- a/numba_cuda/numba/cuda/random.py
+++ b/numba_cuda/numba/cuda/random.py
@@ -3,7 +3,7 @@
 
 import math
 
-from numba import cuda, jit
+from numba import cuda
 from numba.cuda import (
     float32,
     float64,
@@ -11,9 +11,13 @@
     int64,
     uint64,
     from_dtype,
+    HAS_NUMBA,
 )
 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 bcfe09a23..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.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,7 +64,7 @@ def resolve_value_type(self, val):
         from numba.cuda.dispatcher import CUDADispatcher
         from numba.core.dispatcher import Dispatcher
 
-        if _HAS_NUMBA:
+        if HAS_NUMBA:
             if isinstance(val, Dispatcher) and not isinstance(
                 val, CUDADispatcher
             ):
diff --git a/numba_cuda/numba/cuda/testing.py b/numba_cuda/numba/cuda/testing.py
index caf93a831..733cd96fb 100644
--- a/numba_cuda/numba/cuda/testing.py
+++ b/numba_cuda/numba/cuda/testing.py
@@ -18,7 +18,7 @@
 from io import StringIO
 import unittest
 import numpy as np
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
 if PYVERSION >= (3, 10):
     from filecheck.matcher import Matcher
@@ -192,7 +192,7 @@ def skip_on_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)
+    return unittest.skipIf(not HAS_NUMBA, reason)
 
 
 def skip_unless_cudasim(reason):
diff --git a/numba_cuda/numba/cuda/tests/core/test_serialize.py b/numba_cuda/numba/cuda/tests/core/test_serialize.py
index 94683fe56..c4e830faf 100644
--- a/numba_cuda/numba/cuda/tests/core/test_serialize.py
+++ b/numba_cuda/numba/cuda/tests/core/test_serialize.py
@@ -10,10 +10,10 @@
 import unittest
 from multiprocessing import get_context
 
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 import numba
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba.core.errors import TypingError
     from numba.core.target_extension import resolve_dispatcher_from_str
 else:
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 adba8252b..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
@@ -4,7 +4,7 @@
 import sys
 
 from numba import cuda
-from numba.cuda import _HAS_NUMBA
+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
 
@@ -23,7 +23,7 @@ def target_shared_assign(r, x):
 
 assign_cuda_kernel = cuda.jit(cache=True)(target_shared_assign)
 assign_cuda = CUDAUseCase(assign_cuda_kernel)
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba import njit
 
     assign_cpu_jitted = njit(cache=True)(target_shared_assign)
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 b21011dbf..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,9 +5,9 @@
 import itertools
 import numpy as np
 from numba import cuda
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba.core.errors import TypingError
 else:
     from numba.cuda.core.errors import TypingError
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 b2ba8d9a8..3fa62728b 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_device_func.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_device_func.py
@@ -19,9 +19,9 @@
 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
+from numba.cuda import HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba import jit
 
 
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_enums.py b/numba_cuda/numba/cuda/tests/cudapy/test_enums.py
index c5c5d4a2a..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,9 @@
 import numpy as np
 
 from numba.cuda import int16, int32
-from numba.cuda import vectorize, _HAS_NUMBA
+from numba.cuda import vectorize, HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba import njit
 from numba import cuda
 from numba.cuda import types
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_extending.py b/numba_cuda/numba/cuda/tests/cudapy/test_extending.py
index 4c1b56b85..df5afccb1 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_extending.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_extending.py
@@ -8,7 +8,7 @@
 import numpy as np
 import os
 from numba import cuda
-from numba.cuda import _HAS_NUMBA
+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
@@ -32,7 +32,7 @@ def width(self):
         return self.hi - self.lo
 
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba import njit
 else:
     njit = None
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_intrinsics.py b/numba_cuda/numba/cuda/tests/cudapy/test_intrinsics.py
index 9072ff50f..93fbec026 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_intrinsics.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_intrinsics.py
@@ -7,9 +7,9 @@
 import re
 from numba import cuda
 from numba.cuda import int64
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba.core.errors import TypingError
 else:
     from numba.cuda.core.errors import TypingError
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_ir.py b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
index 994a9168c..634ac15ec 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py
@@ -6,9 +6,9 @@
 import warnings
 import numpy as np
 
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba import objmode
     from numba import njit
 from numba.cuda.core import ir
@@ -400,7 +400,7 @@ def foo(a, b, c=12, d=1j, e=None):
                 if np.abs(i) > 0:
                     k = h / i
                     l = np.arange(1, c + 1)
-                    if _HAS_NUMBA:
+                    if HAS_NUMBA:
                         with objmode():
                             print(e, k)
                     m = np.sqrt(l - g)
@@ -416,7 +416,7 @@ def foo(a, b, c=12, d=1j, e=None):
                     for r in range(len(p)):
                         q.append(p[r])
                         if r > 4 + 1:
-                            if _HAS_NUMBA:
+                            if HAS_NUMBA:
                                 with objmode(s="intp", t="complex128"):
                                     s = 123
                                     t = 5
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_localmem.py b/numba_cuda/numba/cuda/tests/cudapy/test_localmem.py
index 25cc4b1ef..8f264be52 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_localmem.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_localmem.py
@@ -6,9 +6,9 @@
 from numba import cuda
 from numba.cuda import int32, complex128, void
 from numba.cuda import types
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba.core.errors import TypingError
 else:
     from numba.cuda.core.errors import TypingError
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_overload.py b/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
index 9fbab1590..2bba8f4e7 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_overload.py
@@ -3,9 +3,9 @@
 
 from numba import cuda
 from numba.cuda import types
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba.core.errors import TypingError
     from numba import njit
     import numba
@@ -366,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 _HAS_NUMBA and numba.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 36d1e64db..659a13cc7 100644
--- a/numba_cuda/numba/cuda/tests/cudapy/test_sm.py
+++ b/numba_cuda/numba/cuda/tests/cudapy/test_sm.py
@@ -3,9 +3,9 @@
 
 from numba import cuda
 from numba.cuda import int32, float64, void
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba.core.errors import TypingError as NumbaTypingError
 from numba.cuda.core.errors import TypingError
 from numba.cuda import types
@@ -413,7 +413,7 @@ def unsupported_type():
         def invalid_string_type():
             arr = cuda.shared.array(10, dtype="int33")  # noqa: F841
 
-        if _HAS_NUMBA:
+        if HAS_NUMBA:
             with self.assertRaisesRegex(NumbaTypingError, rgx):
                 cuda.jit(void())(invalid_string_type)
         else:
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 0fc044b7f..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,9 +4,9 @@
 import numpy as np
 from numba import cuda
 from numba.cuda import float32, int32, void
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba.core.errors import TypingError
 else:
     from numba.cuda.core.errors import TypingError
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_ufuncs.py b/numba_cuda/numba/cuda/tests/cudapy/test_ufuncs.py
index c36101706..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,9 @@
 import warnings
 import numpy as np
 import unittest
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba import njit
 from numba import cuda
 from numba.cuda import config, types
diff --git a/numba_cuda/numba/cuda/tests/support.py b/numba_cuda/numba/cuda/tests/support.py
index e60fb969d..848754e94 100644
--- a/numba_cuda/numba/cuda/tests/support.py
+++ b/numba_cuda/numba/cuda/tests/support.py
@@ -36,9 +36,9 @@
 from numba.cuda.datamodel.models import OpaqueModel
 from numba.cuda.np import numpy_support
 
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba.core.extending import (
         typeof_impl as upstream_typeof_impl,
     )
@@ -770,7 +770,7 @@ def typeof_dummy(val, c):
             return dummy_type
 
         # Dual registration for cross-target tests
-        if _HAS_NUMBA:
+        if HAS_NUMBA:
             UpstreamDummyType = type(
                 "DummyTypeFor{}".format(test_id), (upstream_types.Opaque,), {}
             )
diff --git a/numba_cuda/numba/cuda/tests/test_extending.py b/numba_cuda/numba/cuda/tests/test_extending.py
index 17331693b..9798de584 100644
--- a/numba_cuda/numba/cuda/tests/test_extending.py
+++ b/numba_cuda/numba/cuda/tests/test_extending.py
@@ -11,9 +11,9 @@
 import numba
 from numba.cuda import jit
 from numba.cuda import types
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba.core import errors
 else:
     from numba.cuda.core import errors
diff --git a/numba_cuda/numba/cuda/tests/test_extending_types.py b/numba_cuda/numba/cuda/tests/test_extending_types.py
index cee9222a3..60b465997 100644
--- a/numba_cuda/numba/cuda/tests/test_extending_types.py
+++ b/numba_cuda/numba/cuda/tests/test_extending_types.py
@@ -7,9 +7,9 @@
 
 from numba.cuda import jit
 from numba.cuda import types
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba.core.errors import TypingError, NumbaTypeError
 else:
     from numba.cuda.core.errors import TypingError, NumbaTypeError
diff --git a/numba_cuda/numba/cuda/typing/templates.py b/numba_cuda/numba/cuda/typing/templates.py
index e9078f6e5..d186a0a01 100644
--- a/numba_cuda/numba/cuda/typing/templates.py
+++ b/numba_cuda/numba/cuda/typing/templates.py
@@ -25,9 +25,9 @@
 from numba.cuda import utils
 from numba.cuda.core import targetconfig
 
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 
-if _HAS_NUMBA:
+if HAS_NUMBA:
     from numba.core.typing import Signature as CoreSignature
 
 # info store for inliner callback functions e.g. cost model
@@ -98,7 +98,7 @@ def __hash__(self):
 
     def __eq__(self, other):
         sig_types = (Signature,)
-        if _HAS_NUMBA:
+        if HAS_NUMBA:
             sig_types = (Signature, CoreSignature)
         if isinstance(other, sig_types):
             return (
@@ -382,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 _HAS_NUMBA:
+            if HAS_NUMBA:
                 sig_types = (Signature, CoreSignature)
             if not isinstance(sig, sig_types):
                 raise AssertionError(
diff --git a/numba_cuda/numba/cuda/typing/typeof.py b/numba_cuda/numba/cuda/typing/typeof.py
index caece355c..a3091d282 100644
--- a/numba_cuda/numba/cuda/typing/typeof.py
+++ b/numba_cuda/numba/cuda/typing/typeof.py
@@ -8,7 +8,7 @@
 
 import numpy as np
 from numpy.random.bit_generator import BitGenerator
-from numba.cuda import _HAS_NUMBA
+from numba.cuda import HAS_NUMBA
 from numba.cuda import types
 from numba.cuda.core import errors
 from numba.cuda import utils
@@ -65,7 +65,7 @@ def typeof_impl(val, c):
         if cffi_utils.is_ffi_instance(val):
             return types.ffi
 
-    if _HAS_NUMBA:
+    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
 

From 1ac7c92573025822329ca16a7c2a76f31e6d27d7 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Thu, 30 Oct 2025 12:37:35 -0700
Subject: [PATCH 20/25] remove redundant version_info, simplified code that
 uses it to always assume the numba version > 0.60.0

---
 numba_cuda/_version.py            | 44 -------------------------------
 numba_cuda/numba/cuda/__init__.py |  5 ----
 numba_cuda/numba/cuda/lowering.py | 24 +++++------------
 3 files changed, 7 insertions(+), 66 deletions(-)

diff --git a/numba_cuda/_version.py b/numba_cuda/_version.py
index 83d5a3894..b5622732a 100644
--- a/numba_cuda/_version.py
+++ b/numba_cuda/_version.py
@@ -2,7 +2,6 @@
 # SPDX-License-Identifier: BSD-2-Clause
 
 import importlib.resources
-from collections import namedtuple
 
 __version__ = (
     importlib.resources.files("numba_cuda")
@@ -10,46 +9,3 @@
     .read_text()
     .strip()
 )
-
-version_info = namedtuple(
-    "version_info", ("major minor patch short full string tuple git_revision")
-)
-
-
-def generate_version_info(version):
-    """Process a version string into a structured version_info object.
-
-    Parameters
-    ----------
-    version: str
-        a string describing the current version
-
-    Returns
-    -------
-    version_info: tuple
-        structured version information
-
-    See also
-    --------
-    Look at the definition of 'version_info' in this module for details.
-
-    """
-    parts = version.split(".")
-
-    def try_int(x):
-        try:
-            return int(x)
-        except ValueError:
-            return None
-
-    major = try_int(parts[0]) if len(parts) >= 1 else None
-    minor = try_int(parts[1]) if len(parts) >= 2 else None
-    patch = try_int(parts[2]) if len(parts) >= 3 else None
-    short = (major, minor)
-    full = (major, minor, patch)
-    string = version
-    tup = tuple(parts)
-    git_revision = tup[3] if len(tup) >= 4 else None
-    return version_info(
-        major, minor, patch, short, full, string, tup, git_revision
-    )
diff --git a/numba_cuda/numba/cuda/__init__.py b/numba_cuda/numba/cuda/__init__.py
index 29ac69d68..b56e22eab 100644
--- a/numba_cuda/numba/cuda/__init__.py
+++ b/numba_cuda/numba/cuda/__init__.py
@@ -9,11 +9,6 @@
 
 from numba_cuda._version import __version__
 
-from numba_cuda._version import generate_version_info
-
-version_info = generate_version_info(__version__)
-del generate_version_info
-
 # Re-export types itself
 import numba.cuda.types as types
 
diff --git a/numba_cuda/numba/cuda/lowering.py b/numba_cuda/numba/cuda/lowering.py
index 9df637fd4..0f47ea6b9 100644
--- a/numba_cuda/numba/cuda/lowering.py
+++ b/numba_cuda/numba/cuda/lowering.py
@@ -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.cuda 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:

From 2fbfa1312ca013a639c4500e192f620059cc8047 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Fri, 31 Oct 2025 08:27:00 -0700
Subject: [PATCH 21/25] Revert "Remove MakeFunctionToJitFunction pass, which
 implicitly uses njit, not a use case we intend to support"

This reverts commit 0a60d27f5a0492c30b50393e3fc1fc395b308b35.
---
 numba_cuda/numba/cuda/compiler.py            |  6 ++
 numba_cuda/numba/cuda/core/untyped_passes.py | 67 ++++++++++++++++++++
 2 files changed, 73 insertions(+)

diff --git a/numba_cuda/numba/cuda/compiler.py b/numba_cuda/numba/cuda/compiler.py
index 8d76bc675..a1653c192 100644
--- a/numba_cuda/numba/cuda/compiler.py
+++ b/numba_cuda/numba/cuda/compiler.py
@@ -48,6 +48,7 @@
     WithLifting,
     InlineInlinables,
     FindLiterallyCalls,
+    MakeFunctionToJitFunction,
     LiteralUnroll,
     ReconstructSSA,
     RewriteDynamicRaises,
@@ -224,6 +225,11 @@ def define_untyped_pipeline(state, name="untyped"):
 
         pm.add_pass(RewriteDynamicRaises, "rewrite dynamic raises")
 
+        # convert any remaining closures into functions
+        pm.add_pass(
+            MakeFunctionToJitFunction,
+            "convert make_function into JIT functions",
+        )
         # inline functions that have been determined as inlinable and rerun
         # branch pruning, this needs to be run after closures are inlined as
         # the IR repr of a closure masks call sites if an inlinable is called
diff --git a/numba_cuda/numba/cuda/core/untyped_passes.py b/numba_cuda/numba/cuda/core/untyped_passes.py
index a6f20a2f1..4006854ce 100644
--- a/numba_cuda/numba/cuda/core/untyped_passes.py
+++ b/numba_cuda/numba/cuda/core/untyped_passes.py
@@ -35,6 +35,7 @@
     resolve_func_from_module,
     simplify_CFG,
     GuardException,
+    convert_code_obj_to_function,
     build_definitions,
     replace_var_names,
     get_name_var_table,
@@ -617,6 +618,72 @@ def run_pass(self, state):
         return False
 
 
+@register_pass(mutates_CFG=True, analysis_only=False)
+class MakeFunctionToJitFunction(FunctionPass):
+    """
+    This swaps an ir.Expr.op == "make_function" i.e. a closure, for a compiled
+    function containing the closure body and puts it in ir.Global. It's a 1:1
+    statement value swap. `make_function` is already untyped
+    """
+
+    _name = "make_function_op_code_to_jit_function"
+
+    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():
+            for stmt in blk.body:
+                if isinstance(stmt, ir.Assign):
+                    if isinstance(stmt.value, ir.Expr):
+                        if stmt.value.op == "make_function":
+                            node = stmt.value
+                            getdef = func_ir.get_definition
+                            kw_default = getdef(node.defaults)
+                            ok = False
+                            if kw_default is None or isinstance(
+                                kw_default, ir.Const
+                            ):
+                                ok = True
+                            elif isinstance(kw_default, tuple):
+                                ok = all(
+                                    [
+                                        isinstance(getdef(x), ir.Const)
+                                        for x in kw_default
+                                    ]
+                                )
+                            elif isinstance(kw_default, ir.Expr):
+                                if kw_default.op != "build_tuple":
+                                    continue
+                                ok = all(
+                                    [
+                                        isinstance(getdef(x), ir.Const)
+                                        for x in kw_default.items
+                                    ]
+                                )
+                            if not ok:
+                                continue
+
+                            pyfunc = convert_code_obj_to_function(node, func_ir)
+                            func = njit()(pyfunc)
+                            new_node = ir.Global(
+                                node.code.co_name, func, stmt.loc
+                            )
+                            stmt.value = new_node
+                            mutated |= True
+
+        # if a change was made the del ordering is probably wrong, patch up
+        if mutated:
+            post_proc = postproc.PostProcessor(func_ir)
+            post_proc.run()
+
+        return mutated
+
+
 @register_pass(mutates_CFG=True, analysis_only=False)
 class TransformLiteralUnrollConstListToTuple(FunctionPass):
     """This pass spots a `literal_unroll([<constant values>])` and rewrites it

From 369368b5b9b8e3e7a3e7015b09dc3d09459e31d7 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Fri, 31 Oct 2025 08:30:40 -0700
Subject: [PATCH 22/25] add another dirname for the basepath since numba.cuda
 adds a level in the directory chain

---
 numba_cuda/numba/cuda/typing/templates.py | 24 +++++++++++++++++------
 1 file changed, 18 insertions(+), 6 deletions(-)

diff --git a/numba_cuda/numba/cuda/typing/templates.py b/numba_cuda/numba/cuda/typing/templates.py
index d186a0a01..de4c3b464 100644
--- a/numba_cuda/numba/cuda/typing/templates.py
+++ b/numba_cuda/numba/cuda/typing/templates.py
@@ -407,7 +407,9 @@ def unpack_opt(x):
 
     def get_template_info(self):
         impl = getattr(self, "generic")
-        basepath = os.path.dirname(os.path.dirname(numba.cuda.__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))
@@ -494,7 +496,9 @@ def unpack_opt(x):
 
     def get_template_info(self):
         impl = getattr(self, "generic")
-        basepath = os.path.dirname(os.path.dirname(numba.cuda.__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 = {
@@ -883,7 +887,9 @@ def get_source_info(cls):
             - "docstring": str
                 The docstring of the definition.
         """
-        basepath = os.path.dirname(os.path.dirname(numba.cuda.__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))
@@ -898,7 +904,9 @@ def get_source_info(cls):
         return info
 
     def get_template_info(self):
-        basepath = os.path.dirname(os.path.dirname(numba.cuda.__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))
@@ -1055,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.cuda.__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))
@@ -1217,7 +1227,9 @@ def generic(_, args, kws):
                     return sig.as_method()
 
             def get_template_info(self):
-                basepath = os.path.dirname(os.path.dirname(numba.cuda.__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))

From 184a5294342718bdf4735bb52783731b981cfbd3 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Fri, 31 Oct 2025 08:31:08 -0700
Subject: [PATCH 23/25] remove individual test skips, there's already one for
 the module

---
 .../numba/cuda/tests/core/test_serialize.py      | 16 ----------------
 1 file changed, 16 deletions(-)

diff --git a/numba_cuda/numba/cuda/tests/core/test_serialize.py b/numba_cuda/numba/cuda/tests/core/test_serialize.py
index c4e830faf..d2da82c47 100644
--- a/numba_cuda/numba/cuda/tests/core/test_serialize.py
+++ b/numba_cuda/numba/cuda/tests/core/test_serialize.py
@@ -59,7 +59,6 @@ def check_result(func):
             new_func = pickle.loads(pickled)
             check_result(new_func)
 
-    @skip_on_standalone_numba_cuda
     def test_call_with_sig(self):
         from .serialize_usecases import add_with_sig
 
@@ -67,7 +66,6 @@ def test_call_with_sig(self):
         # Compilation has been disabled => float inputs will be coerced to int
         self.run_with_protocols(self.check_call, add_with_sig, 5, (1.2, 4.2))
 
-    @skip_on_standalone_numba_cuda
     def test_call_without_sig(self):
         from .serialize_usecases import add_without_sig
 
@@ -80,7 +78,6 @@ def test_call_without_sig(self):
             self.check_call, add_without_sig, "abc", ("a", "bc")
         )
 
-    @skip_on_standalone_numba_cuda
     def test_call_nopython(self):
         from .serialize_usecases import add_nopython
 
@@ -90,7 +87,6 @@ def test_call_nopython(self):
             self.check_call, add_nopython, TypingError, (object(), object())
         )
 
-    @skip_on_standalone_numba_cuda
     def test_call_nopython_fail(self):
         from .serialize_usecases import add_nopython_fail
 
@@ -99,7 +95,6 @@ def test_call_nopython_fail(self):
             self.check_call, add_nopython_fail, TypingError, (1, 2)
         )
 
-    @skip_on_standalone_numba_cuda
     def test_call_objmode_with_global(self):
         from .serialize_usecases import get_global_objmode
 
@@ -107,14 +102,12 @@ def test_call_objmode_with_global(self):
             self.check_call, get_global_objmode, 7.5, (2.5,)
         )
 
-    @skip_on_standalone_numba_cuda
     def test_call_closure(self):
         from .serialize_usecases import closure
 
         inner = closure(1)
         self.run_with_protocols(self.check_call, inner, 6, (2, 3))
 
-    @skip_on_standalone_numba_cuda
     def check_call_closure_with_globals(self, **jit_args):
         from .serialize_usecases import closure_with_globals
 
@@ -127,35 +120,30 @@ def test_call_closure_with_globals_nopython(self):
     def test_call_closure_with_globals_objmode(self):
         self.check_call_closure_with_globals(forceobj=True)
 
-    @skip_on_standalone_numba_cuda
     def test_call_closure_calling_other_function(self):
         from .serialize_usecases import closure_calling_other_function
 
         inner = closure_calling_other_function(3.0)
         self.run_with_protocols(self.check_call, inner, 11.0, (4.0, 6.0))
 
-    @skip_on_standalone_numba_cuda
     def test_call_closure_calling_other_closure(self):
         from .serialize_usecases import closure_calling_other_closure
 
         inner = closure_calling_other_closure(3.0)
         self.run_with_protocols(self.check_call, inner, 8.0, (4.0,))
 
-    @skip_on_standalone_numba_cuda
     def test_call_dyn_func(self):
         from .serialize_usecases import dyn_func
 
         # Check serializing a dynamically-created function
         self.run_with_protocols(self.check_call, dyn_func, 36, (6,))
 
-    @skip_on_standalone_numba_cuda
     def test_call_dyn_func_objmode(self):
         from .serialize_usecases import dyn_func_objmode
 
         # Same with an object mode function
         self.run_with_protocols(self.check_call, dyn_func_objmode, 36, (6,))
 
-    @skip_on_standalone_numba_cuda
     def test_renamed_module(self):
         from .serialize_usecases import get_renamed_module
 
@@ -166,7 +154,6 @@ def test_renamed_module(self):
             self.check_call, get_renamed_module, expected, (0.0,)
         )
 
-    @skip_on_standalone_numba_cuda
     def test_other_process(self):
         """
         Check that reconstructing doesn't depend on resources already
@@ -186,7 +173,6 @@ def test_other_process(self):
             """.format(**locals())
         subprocess.check_call([sys.executable, "-c", code])
 
-    @skip_on_standalone_numba_cuda
     def test_reuse(self):
         """
         Check that deserializing the same function multiple times re-uses
@@ -226,7 +212,6 @@ def test_reuse(self):
         g.disable_compile()
         self.assertEqual(g(2, 4), 13)
 
-    @skip_on_standalone_numba_cuda
     def test_imp_deprecation(self):
         """
         The imp module was deprecated in v3.4 in favour of importlib
@@ -318,7 +303,6 @@ class Klass:
         proc.join(timeout=60)
         self.assertEqual(proc.exitcode, 0)
 
-    @skip_on_standalone_numba_cuda
     def test_dynamic_class_issue_7356(self):
         import numba
 

From cc4682a3560302ea0baafa3eb2a0df15180ab4b6 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Fri, 31 Oct 2025 08:46:49 -0700
Subject: [PATCH 24/25] address more review comments

---
 numba_cuda/numba/cuda/__init__.py             |   4 -
 numba_cuda/numba/cuda/core/ir.py              |   3 +-
 numba_cuda/numba/cuda/core/untyped_passes.py  |   5 +-
 numba_cuda/numba/cuda/cpython/listobj.py      |  43 +--
 numba_cuda/numba/cuda/misc/mergesort.py       | 124 --------
 numba_cuda/numba/cuda/misc/quicksort.py       | 270 ------------------
 numba_cuda/numba/cuda/np/arraymath.py         |  69 -----
 numba_cuda/numba/cuda/np/arrayobj.py          |  31 +-
 .../np/polynomial/polynomial_functions.py     |   2 +-
 numba_cuda/numba/cuda/random.py               |   2 +-
 numba_cuda/numba/cuda/typing/arraydecl.py     |  21 --
 numba_cuda/numba/cuda/typing/npydecl.py       |   6 -
 12 files changed, 8 insertions(+), 572 deletions(-)
 delete mode 100644 numba_cuda/numba/cuda/misc/mergesort.py
 delete mode 100644 numba_cuda/numba/cuda/misc/quicksort.py

diff --git a/numba_cuda/numba/cuda/__init__.py b/numba_cuda/numba/cuda/__init__.py
index b56e22eab..d0ff4ba55 100644
--- a/numba_cuda/numba/cuda/__init__.py
+++ b/numba_cuda/numba/cuda/__init__.py
@@ -7,7 +7,6 @@
 import warnings
 import sys
 
-from numba_cuda._version import __version__
 
 # Re-export types itself
 import numba.cuda.types as types
@@ -69,6 +68,3 @@
     )
 
 from numba.cuda.np.ufunc import vectorize, guvectorize
-from numba.cuda.misc import quicksort, mergesort
-from numba.cuda.misc.special import literal_unroll
-from numba.cuda.np.numpy_support import *
diff --git a/numba_cuda/numba/cuda/core/ir.py b/numba_cuda/numba/cuda/core/ir.py
index b907b2f14..807733fea 100644
--- a/numba_cuda/numba/cuda/core/ir.py
+++ b/numba_cuda/numba/cuda/core/ir.py
@@ -28,7 +28,6 @@
     ConstantInferenceError,
 )
 from numba.cuda.core import consts
-from numba.cuda import HAS_NUMBA
 
 # terminal color markup
 _termcolor = errors.termcolor()
@@ -783,7 +782,7 @@ def __init__(self, value, loc):
         if HAS_NUMBA:
             assert isinstance(loc, (Loc, numba.core.ir.Loc))
         else:
-            assert isinstance(loc)
+            assert isinstance(loc, (Loc))
         self.value = value
         self.loc = loc
 
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/listobj.py b/numba_cuda/numba/cuda/cpython/listobj.py
index de28b6399..ee059a8e7 100644
--- a/numba_cuda/numba/cuda/cpython/listobj.py
+++ b/numba_cuda/numba/cuda/cpython/listobj.py
@@ -19,9 +19,8 @@
     RefType,
 )
 from numba.cuda.extending import overload_method, overload
-from numba.cuda.misc import quicksort
 from numba.cuda.cpython import slicing
-from numba.cuda 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/misc/mergesort.py b/numba_cuda/numba/cuda/misc/mergesort.py
deleted file mode 100644
index c6ed91993..000000000
--- a/numba_cuda/numba/cuda/misc/mergesort.py
+++ /dev/null
@@ -1,124 +0,0 @@
-# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
-# SPDX-License-Identifier: BSD-2-Clause
-
-"""
-The same algorithm as translated from numpy.
-See numpy/core/src/npysort/mergesort.c.src.
-The high-level numba code is adding a little overhead comparing to
-the pure-C implementation in numpy.
-"""
-
-import numpy as np
-from collections import namedtuple
-
-# Array size smaller than this will be sorted by insertion sort
-SMALL_MERGESORT = 20
-
-
-MergesortImplementation = namedtuple(
-    "MergesortImplementation",
-    [
-        "run_mergesort",
-    ],
-)
-
-
-def make_mergesort_impl(lt=None, is_argsort=False):
-    # The less than
-    if lt is None:
-
-        def lt(a, b):
-            return a < b
-    else:
-        lt = lt
-
-    if is_argsort:
-
-        def lessthan(a, b, vals):
-            return lt(vals[a], vals[b])
-    else:
-
-        def lessthan(a, b, vals):
-            return lt(a, b)
-
-    def argmergesort_inner(arr, vals, ws):
-        """The actual mergesort function
-
-        Parameters
-        ----------
-        arr : array [read+write]
-            The values being sorted inplace.  For argsort, this is the
-            indices.
-        vals : array [readonly]
-            ``None`` for normal sort.  In argsort, this is the actual array values.
-        ws : array [write]
-            The workspace.  Must be of size ``arr.size // 2``
-        """
-        if arr.size > SMALL_MERGESORT:
-            # Merge sort
-            mid = arr.size // 2
-
-            argmergesort_inner(arr[:mid], vals, ws)
-            argmergesort_inner(arr[mid:], vals, ws)
-
-            # Copy left half into workspace so we don't overwrite it
-            for i in range(mid):
-                ws[i] = arr[i]
-
-            # Merge
-            left = ws[:mid]
-            right = arr[mid:]
-            out = arr
-
-            i = j = k = 0
-            while i < left.size and j < right.size:
-                if not lessthan(right[j], left[i], vals):
-                    out[k] = left[i]
-                    i += 1
-                else:
-                    out[k] = right[j]
-                    j += 1
-                k += 1
-
-            # Leftovers
-            while i < left.size:
-                out[k] = left[i]
-                i += 1
-                k += 1
-
-            while j < right.size:
-                out[k] = right[j]
-                j += 1
-                k += 1
-        else:
-            # Insertion sort
-            i = 1
-            while i < arr.size:
-                j = i
-                while j > 0 and lessthan(arr[j], arr[j - 1], vals):
-                    arr[j - 1], arr[j] = arr[j], arr[j - 1]
-                    j -= 1
-                i += 1
-
-    # The top-level entry points
-
-    def mergesort(arr):
-        "Inplace"
-        ws = np.empty(arr.size // 2, dtype=arr.dtype)
-        argmergesort_inner(arr, None, ws)
-        return arr
-
-    def argmergesort(arr):
-        "Out-of-place"
-        idxs = np.arange(arr.size)
-        ws = np.empty(arr.size // 2, dtype=idxs.dtype)
-        argmergesort_inner(idxs, arr, ws)
-        return idxs
-
-    return MergesortImplementation(
-        run_mergesort=(argmergesort if is_argsort else mergesort)
-    )
-
-
-def make_jit_mergesort(*args, **kwargs):
-    return make_mergesort_impl(*args, **kwargs)
diff --git a/numba_cuda/numba/cuda/misc/quicksort.py b/numba_cuda/numba/cuda/misc/quicksort.py
deleted file mode 100644
index 506e3636d..000000000
--- a/numba_cuda/numba/cuda/misc/quicksort.py
+++ /dev/null
@@ -1,270 +0,0 @@
-# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
-# SPDX-License-Identifier: BSD-2-Clause
-
-import collections
-
-import numpy as np
-
-from numba.cuda import types
-
-
-QuicksortImplementation = collections.namedtuple(
-    "QuicksortImplementation",
-    (  # The compile function itself
-        "compile",
-        # All subroutines exercised by test_sort
-        "partition",
-        "partition3",
-        "insertion_sort",
-        # The top-level function
-        "run_quicksort",
-    ),
-)
-
-
-Partition = collections.namedtuple("Partition", ("start", "stop"))
-
-# Under this size, switch to a simple insertion sort
-SMALL_QUICKSORT = 15
-
-MAX_STACK = 100
-
-
-def make_quicksort_impl(
-    wrap, lt=None, is_argsort=False, is_list=False, is_np_array=False
-):
-    intp = types.intp
-    zero = intp(0)
-
-    # Two subroutines to make the core algorithm generic wrt. argsort
-    # or normal sorting.  Note the genericity may make basic sort()
-    # slightly slower (~5%)
-    if is_argsort:
-        if is_list:
-
-            @wrap
-            def make_res(A):
-                return [x for x in range(len(A))]
-        else:
-
-            @wrap
-            def make_res(A):
-                return np.arange(A.size)
-
-        @wrap
-        def GET(A, idx_or_val):
-            return A[idx_or_val]
-
-    else:
-
-        @wrap
-        def make_res(A):
-            return A
-
-        @wrap
-        def GET(A, idx_or_val):
-            return idx_or_val
-
-    def default_lt(a, b):
-        """
-        Trivial comparison function between two keys.
-        """
-        return a < b
-
-    LT = wrap(lt if lt is not None else default_lt)
-
-    @wrap
-    def insertion_sort(A, R, low, high):
-        """
-        Insertion sort A[low:high + 1]. Note the inclusive bounds.
-        """
-        assert low >= 0
-        if high <= low:
-            return
-
-        for i in range(low + 1, high + 1):
-            k = R[i]
-            v = GET(A, k)
-            # Insert v into A[low:i]
-            j = i
-            while j > low and LT(v, GET(A, R[j - 1])):
-                # Make place for moving A[i] downwards
-                R[j] = R[j - 1]
-                j -= 1
-            R[j] = k
-
-    @wrap
-    def partition(A, R, low, high):
-        """
-        Partition A[low:high + 1] around a chosen pivot.  The pivot's index
-        is returned.
-        """
-        assert low >= 0
-        assert high > low
-
-        mid = (low + high) >> 1
-        # NOTE: the pattern of swaps below for the pivot choice and the
-        # partitioning gives good results (i.e. regular O(n log n))
-        # on sorted, reverse-sorted, and uniform arrays.  Subtle changes
-        # risk breaking this property.
-
-        # median of three {low, middle, high}
-        if LT(GET(A, R[mid]), GET(A, R[low])):
-            R[low], R[mid] = R[mid], R[low]
-        if LT(GET(A, R[high]), GET(A, R[mid])):
-            R[high], R[mid] = R[mid], R[high]
-        if LT(GET(A, R[mid]), GET(A, R[low])):
-            R[low], R[mid] = R[mid], R[low]
-        pivot = GET(A, R[mid])
-
-        # Temporarily stash the pivot at the end
-        R[high], R[mid] = R[mid], R[high]
-        i = low
-        j = high - 1
-        while True:
-            while i < high and LT(GET(A, R[i]), pivot):
-                i += 1
-            while j >= low and LT(pivot, GET(A, R[j])):
-                j -= 1
-            if i >= j:
-                break
-            R[i], R[j] = R[j], R[i]
-            i += 1
-            j -= 1
-        # Put the pivot back in its final place (all items before `i`
-        # are smaller than the pivot, all items at/after `i` are larger)
-        R[i], R[high] = R[high], R[i]
-        return i
-
-    @wrap
-    def partition3(A, low, high):
-        """
-        Three-way partition [low, high) around a chosen pivot.
-        A tuple (lt, gt) is returned such that:
-            - all elements in [low, lt) are < pivot
-            - all elements in [lt, gt] are == pivot
-            - all elements in (gt, high] are > pivot
-        """
-        mid = (low + high) >> 1
-        # median of three {low, middle, high}
-        if LT(A[mid], A[low]):
-            A[low], A[mid] = A[mid], A[low]
-        if LT(A[high], A[mid]):
-            A[high], A[mid] = A[mid], A[high]
-        if LT(A[mid], A[low]):
-            A[low], A[mid] = A[mid], A[low]
-        pivot = A[mid]
-
-        A[low], A[mid] = A[mid], A[low]
-        lt = low
-        gt = high
-        i = low + 1
-        while i <= gt:
-            if LT(A[i], pivot):
-                A[lt], A[i] = A[i], A[lt]
-                lt += 1
-                i += 1
-            elif LT(pivot, A[i]):
-                A[gt], A[i] = A[i], A[gt]
-                gt -= 1
-            else:
-                i += 1
-        return lt, gt
-
-    @wrap
-    def run_quicksort1(A):
-        R = make_res(A)
-
-        if len(A) < 2:
-            return R
-
-        stack = [Partition(zero, zero)] * MAX_STACK
-        stack[0] = Partition(zero, len(A) - 1)
-        n = 1
-
-        while n > 0:
-            n -= 1
-            low, high = stack[n]
-            # Partition until it becomes more efficient to do an insertion sort
-            while high - low >= SMALL_QUICKSORT:
-                assert n < MAX_STACK
-                i = partition(A, R, low, high)
-                # Push largest partition on the stack
-                if high - i > i - low:
-                    # Right is larger
-                    if high > i:
-                        stack[n] = Partition(i + 1, high)
-                        n += 1
-                    high = i - 1
-                else:
-                    if i > low:
-                        stack[n] = Partition(low, i - 1)
-                        n += 1
-                    low = i + 1
-
-            insertion_sort(A, R, low, high)
-
-        return R
-
-    if is_np_array:
-
-        @wrap
-        def run_quicksort(A):
-            if A.ndim == 1:
-                return run_quicksort1(A)
-            else:
-                for idx in np.ndindex(A.shape[:-1]):
-                    run_quicksort1(A[idx])
-                return A
-    else:
-
-        @wrap
-        def run_quicksort(A):
-            return run_quicksort1(A)
-
-    # Unused quicksort implementation based on 3-way partitioning; the
-    # partitioning scheme turns out exhibiting bad behaviour on sorted arrays.
-    @wrap
-    def _run_quicksort(A):
-        stack = [Partition(zero, zero)] * 100
-        stack[0] = Partition(zero, len(A) - 1)
-        n = 1
-
-        while n > 0:
-            n -= 1
-            low, high = stack[n]
-            # Partition until it becomes more efficient to do an insertion sort
-            while high - low >= SMALL_QUICKSORT:
-                assert n < MAX_STACK
-                l, r = partition3(A, low, high)
-                # One trivial (empty) partition => iterate on the other
-                if r == high:
-                    high = l - 1
-                elif l == low:
-                    low = r + 1
-                # Push largest partition on the stack
-                elif high - r > l - low:
-                    # Right is larger
-                    stack[n] = Partition(r + 1, high)
-                    n += 1
-                    high = l - 1
-                else:
-                    stack[n] = Partition(low, l - 1)
-                    n += 1
-                    low = r + 1
-
-            insertion_sort(A, low, high)
-
-    return QuicksortImplementation(
-        wrap, partition, partition3, insertion_sort, run_quicksort
-    )
-
-
-def make_py_quicksort(*args, **kwargs):
-    return make_quicksort_impl((lambda f: f), *args, **kwargs)
-
-
-def make_jit_quicksort(*args, **kwargs):
-    from numba.cuda.extending import register_jitable
-
-    return make_quicksort_impl((lambda f: register_jitable(f)), *args, **kwargs)
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 351959ce2..de8834e91 100644
--- a/numba_cuda/numba/cuda/np/arrayobj.py
+++ b/numba_cuda/numba/cuda/np/arrayobj.py
@@ -16,7 +16,7 @@
 
 import numpy as np
 
-from numba.cuda import literal_unroll
+from numba.cuda.misc.special import literal_unroll
 from numba.cuda import types, typing
 from numba.cuda.core import errors
 from numba.cuda import cgutils, extending
@@ -54,7 +54,6 @@
     intrinsic,
     overload_attribute,
 )
-from numba.cuda.misc import quicksort, mergesort
 from numba.cuda.cpython import slicing
 from numba.cuda.cpython.unsafe.tuple import (
     tuple_setitem,
@@ -7125,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):
@@ -7148,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/polynomial/polynomial_functions.py b/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py
index d64e9c104..5d974f8da 100644
--- a/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py
+++ b/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py
@@ -9,7 +9,7 @@
 from numpy.polynomial import polynomial as poly
 from numpy.polynomial import polyutils as pu
 
-from numba.cuda import literal_unroll
+from numba.cuda.misc.special import literal_unroll
 from numba.cuda import types
 from numba.cuda.core import errors
 from numba.cuda.extending import overload
diff --git a/numba_cuda/numba/cuda/random.py b/numba_cuda/numba/cuda/random.py
index 8cf7cc079..e8e620d5c 100644
--- a/numba_cuda/numba/cuda/random.py
+++ b/numba_cuda/numba/cuda/random.py
@@ -10,9 +10,9 @@
     uint32,
     int64,
     uint64,
-    from_dtype,
     HAS_NUMBA,
 )
+from numba.cuda.np.numpy_support import from_dtype
 from numba.cuda import config
 
 if HAS_NUMBA:
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/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))

From 19cc76bad0c05f249b60ec75a8884c0debf352c9 Mon Sep 17 00:00:00 2001
From: Atmn Patel <atmnp@nvidia.com>
Date: Fri, 31 Oct 2025 11:48:03 -0700
Subject: [PATCH 25/25] add guard for JIT_COVERAGE attr in config used

---
 numba_cuda/numba/cuda/misc/coverage_support.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/numba_cuda/numba/cuda/misc/coverage_support.py b/numba_cuda/numba/cuda/misc/coverage_support.py
index 5b50c9c85..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(