merge/resolve

Author: brandon-b-miller

ci/test_thirdparty.sh

@@ -41,13 +41,13 @@ rapids-logger "Show Numba system info"
 python -m numba --sysinfo
 
 rapids-logger "Run Scalar UDF tests"
-python -m pytest python/cudf/cudf/tests/dataframe/methods/test_apply.py -W ignore::UserWarning -W ignore::DeprecationWarning:numba.cuda.core.config
+python -m pytest python/cudf/cudf/tests/dataframe/methods/test_apply.py -W ignore::UserWarning
 
 rapids-logger "Run GroupBy UDF tests"
-python -m pytest python/cudf/cudf/tests/groupby/test_apply.py -k test_groupby_apply_jit -W ignore::UserWarning -W ignore::DeprecationWarning:numba.cuda.core.config
+python -m pytest python/cudf/cudf/tests/groupby/test_apply.py -k test_groupby_apply_jit -W ignore::UserWarning
 
 rapids-logger "Run NRT Stats Counting tests"
-python -m pytest python/cudf/cudf/tests/private_objects/test_nrt_stats.py  -W ignore::UserWarning -W ignore::DeprecationWarning:numba.cuda.core.config
+python -m pytest python/cudf/cudf/tests/private_objects/test_nrt_stats.py  -W ignore::UserWarning
 
 
 popd

numba_cuda/numba/cuda/_internal/cuda_bf16.py

@@ -27,8 +27,8 @@
 )
 from numba.cuda.core.imputils import Registry as TargetRegistry
 from numba.cuda.core.imputils import lower_cast
-from numba.core.typing import signature
-from numba.core.typing.builtins import (
+from numba.cuda.typing import signature
+from numba.cuda.typing.builtins import (
     BinOp,
     BinOpTrueDiv,
     UnaryNegate,

numba_cuda/numba/cuda/_internal/cuda_fp16.py

@@ -29,7 +29,7 @@
 )
 from numba.cuda.core.imputils import Registry as TargetRegistry
 from numba.cuda.core.imputils import lower_cast
-from numba.core.typing import signature
+from numba.cuda.typing import signature
 from numba.cuda import CUSource, declare_device
 from numba.cuda._internal.cuda_bf16 import _type___nv_bfloat16
 from numba.cuda.typing.templates import (

numba_cuda/numba/cuda/api.py

@@ -271,12 +271,8 @@ def open_ipc_array(handle, shape, dtype, strides=None, offset=0):
     # compute size
     size = np.prod(shape) * dtype.itemsize
     # manually recreate the IPC mem handle
-    if driver.USE_NV_BINDING:
-        driver_handle = driver.binding.CUipcMemHandle()
-        driver_handle.reserved = handle
-    else:
-        driver_handle = driver.drvapi.cu_ipc_mem_handle()
-        driver_handle.reserved[:] = handle
+    driver_handle = driver.binding.CUipcMemHandle()
+    driver_handle.reserved = handle
     # use *IpcHandle* to open the IPC memory
     ipchandle = driver.IpcHandle(None, driver_handle, size, offset=offset)
     yield ipchandle.open_array(

numba_cuda/numba/cuda/args.py

@@ -8,7 +8,7 @@
 
 import abc
 
-from numba.core.typing.typeof import typeof, Purpose
+from numba.cuda.typing.typeof import typeof, Purpose
 
 
 class ArgHint(metaclass=abc.ABCMeta):

numba_cuda/numba/cuda/cg.py

@@ -3,7 +3,7 @@
 
 from numba.core import types
 from numba.cuda.extending import overload, overload_method
-from numba.core.typing import signature
+from numba.cuda.typing import signature
 from numba.cuda import nvvmutils
 from numba.cuda.extending import intrinsic
 from numba.cuda.types import grid_group, GridGroup as GridGroupClass

numba_cuda/numba/cuda/core/base.py

@@ -14,9 +14,8 @@
 from numba.core import (
     types,
     datamodel,
-    config,
 )
-from numba.cuda import cgutils, debuginfo, utils
+from numba.cuda import cgutils, debuginfo, utils, config
 from numba.core import errors
 from numba.cuda.core import targetconfig, funcdesc, imputils
 from numba.core.compiler_lock import global_compiler_lock
@@ -357,6 +356,73 @@ def install_registry(self, registry):
         self._insert_cast_defn(loader.new_registrations("casts"))
         self._insert_get_constant_defn(loader.new_registrations("constants"))
 
+    def install_external_registry(self, registry):
+        """
+        Install only third-party registrations from a shared registry like Numba's builtin_registry.
+        Exclude Numba's own implementations in this case (i.e., anything from numba.* namespace).
+
+        This is useful for selectively installing third-party implementations
+        present in the shared builtin_registry from Numba without pulling in any CPU-specific
+        implementations from Numba.
+
+        Note: For getattrs/setattrs, we check the TYPE's __module__ (from the signature)
+        rather than the implementation's __module__, because @lower_getattr/@lower_setattr decorators
+        always set impl.__module__ = "numba.*" regardless of where they are called from.
+        """
+
+        def is_external(obj):
+            """Check if object is from outside numba.* namespace."""
+            try:
+                return not obj.__module__.startswith("numba.")
+            except AttributeError:
+                return True
+
+        def is_external_type_sig(sig):
+            """Check if type in signature is from outside numba.* namespace."""
+            try:
+                return sig and is_external(sig[0])
+            except (AttributeError, IndexError):
+                return True
+
+        try:
+            loader = self._registries[registry]
+        except KeyError:
+            loader = RegistryLoader(registry)
+            self._registries[registry] = loader
+
+        # Filter registrations
+        funcs = [
+            (impl, func, sig)
+            for impl, func, sig in loader.new_registrations("functions")
+            if is_external(impl)
+        ]
+        getattrs = [
+            (impl, attr, sig)
+            for impl, attr, sig in loader.new_registrations("getattrs")
+            if is_external_type_sig(sig)
+        ]
+        setattrs = [
+            (impl, attr, sig)
+            for impl, attr, sig in loader.new_registrations("setattrs")
+            if is_external_type_sig(sig)
+        ]
+        casts = [
+            (impl, sig)
+            for impl, sig in loader.new_registrations("casts")
+            if is_external(impl)
+        ]
+        constants = [
+            (impl, sig)
+            for impl, sig in loader.new_registrations("constants")
+            if is_external(impl)
+        ]
+
+        self.insert_func_defn(funcs)
+        self._insert_getattr_defn(getattrs)
+        self._insert_setattr_defn(setattrs)
+        self._insert_cast_defn(casts)
+        self._insert_get_constant_defn(constants)
+
     def insert_func_defn(self, defns):
         for impl, func, sig in defns:
             self._defns[func].append(impl, sig)
@@ -822,14 +888,15 @@ def _compile_subroutine_no_cache(
         Note this context's flags are not inherited.
         """
         # Compile
-        from numba.core import compiler
+        from numba.cuda import compiler
+        from numba.cuda.flags import Flags
 
         with global_compiler_lock:
             codegen = self.codegen()
             library = codegen.create_library(impl.__name__)
             if flags is None:
                 cstk = targetconfig.ConfigStack()
-                flags = compiler.Flags()
+                flags = Flags()
                 if cstk:
                     tls_flags = cstk.top()
                     if tls_flags.is_set("nrt") and tls_flags.nrt:

numba_cuda/numba/cuda/core/boxing.py

@@ -7,10 +7,11 @@
 
 from llvmlite import ir
 
-from numba.core import types, cgutils
+from numba.core import types
+from numba.cuda import cgutils
 from numba.cuda.core.pythonapi import box, unbox, reflect, NativeValue
 from numba.core.errors import NumbaNotImplementedError, TypingError
-from numba.core.typing.typeof import typeof, Purpose
+from numba.cuda.typing.typeof import typeof, Purpose
 
 from numba.cuda.cpython import setobj, listobj
 from numba.cuda.np import numpy_support

numba_cuda/numba/cuda/core/config.py

@@ -124,49 +124,6 @@ def _process_opt_level(opt_level):
         return _OptLevel(opt_level)
 
 
-class _EnvVar(object):
-    """Descriptor for configuration values that checks numba.config on access."""
-
-    def __init__(self, value, name):
-        self.name = name
-        if isinstance(value, _EnvVar):
-            self.value = value.__get__()
-        else:
-            self.value = value
-        self.check_numba_config()
-
-    def check_numba_config(self):
-        """Check for conflicting value in numba.config and emit deprecation warning."""
-        try:
-            from numba import config as numba_config
-
-            if hasattr(numba_config, self.name):
-                config_value = getattr(numba_config, self.name)
-                if config_value != self.value:
-                    msg = (
-                        f"Configuration value '{self.name}' is explicitly set "
-                        f"to `{config_value}` in numba.config. "
-                        "numba.config is deprecated for numba-cuda "
-                        "and support for configuration values from it "
-                        "will be removed in a future release. "
-                        "Please use numba.cuda.core.config."
-                    )
-                    warnings.warn(msg, category=DeprecationWarning)
-                    self.value = config_value
-            else:
-                # Initialize any missing variables in numba.config
-                setattr(numba_config, self.name, self.value)
-        except ImportError:
-            pass
-
-    def __get__(self):
-        self.check_numba_config()
-        return self.value
-
-    def __set__(self, value):
-        self.value = value
-
-
 class _EnvReloader(object):
     def __init__(self):
         self.reset()
@@ -211,18 +168,7 @@ def update(self, force=False):
         self.validate()
 
     def validate(self):
-        current_module = sys.modules[__name__]
-        try:
-            CUDA_USE_NVIDIA_BINDING = current_module.CUDA_USE_NVIDIA_BINDING
-        except AttributeError:
-            CUDA_USE_NVIDIA_BINDING = 0
-
-        try:
-            CUDA_PER_THREAD_DEFAULT_STREAM = (
-                current_module.CUDA_PER_THREAD_DEFAULT_STREAM
-            )
-        except AttributeError:
-            CUDA_PER_THREAD_DEFAULT_STREAM = 0
+        global CUDA_USE_NVIDIA_BINDING
 
         if CUDA_USE_NVIDIA_BINDING:  # noqa: F821
             try:
@@ -235,7 +181,7 @@ def validate(self):
                 )
                 warnings.warn(msg)
 
-                current_module.CUDA_USE_NVIDIA_BINDING = 0
+                CUDA_USE_NVIDIA_BINDING = False
 
             if CUDA_PER_THREAD_DEFAULT_STREAM:  # noqa: F821
                 warnings.warn(
@@ -250,23 +196,18 @@ def process_environ(self, environ):
         def _readenv(name, ctor, default):
             value = environ.get(name)
             if value is None:
-                result = default() if callable(default) else default
-            else:
-                try:
-                    result = ctor(value)
-                except Exception:
-                    warnings.warn(
-                        f"Environment variable '{name}' is defined but "
-                        f"its associated value '{value}' could not be "
-                        "parsed.\nThe parse failed with exception:\n"
-                        f"{traceback.format_exc()}",
-                        RuntimeWarning,
-                    )
-                    result = default() if callable(default) else default
-            var_name = name
-            if name.startswith("NUMBA_"):
-                var_name = name[6:]
-            return _EnvVar(result, var_name)
+                return default() if callable(default) else default
+            try:
+                return ctor(value)
+            except Exception:
+                warnings.warn(
+                    f"Environment variable '{name}' is defined but "
+                    f"its associated value '{value}' could not be "
+                    "parsed.\nThe parse failed with exception:\n"
+                    f"{traceback.format_exc()}",
+                    RuntimeWarning,
+                )
+                return default
 
         def optional_str(x):
             return str(x) if x is not None else None
@@ -348,12 +289,6 @@ def optional_str(x):
         # Enable NRT statistics counters
         NRT_STATS = _readenv("NUMBA_NRT_STATS", int, 0)
 
-        # Enable NRT statistics
-        CUDA_NRT_STATS = _readenv("NUMBA_CUDA_NRT_STATS", int, 0)
-
-        # Enable NRT
-        CUDA_ENABLE_NRT = _readenv("NUMBA_CUDA_ENABLE_NRT", int, 0)
-
         # How many recently deserialized functions to retain regardless
         # of external references
         FUNCTION_CACHE_SIZE = _readenv("NUMBA_FUNCTION_CACHE_SIZE", int, 128)
@@ -695,53 +630,26 @@ def which_gdb(path_or_bin):
             0,
         )
 
-        # Inject the configuration values into _descriptors
-        if not hasattr(self, "_descriptors"):
-            self._descriptors = {}
-
+        # Inject the configuration values into the module globals
         for name, value in locals().copy().items():
             if name.isupper():
-                self._descriptors[name] = value
+                globals()[name] = value
 
 
 _env_reloader = _EnvReloader()
 
 
-def __getattr__(name):
-    """Module-level __getattr__ provides dynamic behavior for _EnvVar descriptors."""
-    # Fetch non-descriptor globals directly
-    if name in globals():
-        return globals()[name]
-
-    if (
-        hasattr(_env_reloader, "_descriptors")
-        and name in _env_reloader._descriptors
-    ):
-        return _env_reloader._descriptors[name].__get__()
-
-    raise AttributeError(f"module {__name__} has no attribute {name}")
-
-
-def __setattr__(name, value):
-    """Module-level __setattr__ provides dynamic behavior for _EnvVar descriptors."""
-    # Update non-descriptor globals
-    if name in globals():
-        globals()[name] = value
-        return
-
-    if (
-        hasattr(_env_reloader, "_descriptors")
-        and name in _env_reloader._descriptors
-    ):
-        _env_reloader._descriptors[name].__set__(value)
-    else:
-        if not hasattr(_env_reloader, "_descriptors"):
-            _env_reloader._descriptors = {}
-        _env_reloader._descriptors[name] = _EnvVar(value, name)
-
-
 def reload_config():
     """
     Reload the configuration from environment variables, if necessary.
     """
     _env_reloader.update()
+
+
+# use numba.core.config if available, otherwise use numba.cuda.core.config
+try:
+    import numba.core.config as _config
+
+    sys.modules[__name__] = _config
+except ImportError:
+    pass