Merge branch 'main' into vk/typing

Author: VijayKandiah

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/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

numba_cuda/numba/cuda/memory_management/nrt.py

@@ -19,7 +19,7 @@
 )
 from numba.cuda.cudadrv import devices
 from numba.cuda.api import get_current_device
-from numba.cuda.utils import cached_file_read
+from numba.cuda.utils import _readenv, cached_file_read
 from numba.cuda.cudadrv.linkable_code import CUSource
 from numba.cuda.typing.templates import signature
 
@@ -28,6 +28,22 @@
 _nrt_mstats = namedtuple("nrt_mstats", ["alloc", "free", "mi_alloc", "mi_free"])
 
 
+# Check environment variable or config for NRT statistics enablement
+NRT_STATS = _readenv("NUMBA_CUDA_NRT_STATS", bool, False) or getattr(
+    config, "NUMBA_CUDA_NRT_STATS", False
+)
+if not hasattr(config, "NUMBA_CUDA_NRT_STATS"):
+    config.CUDA_NRT_STATS = NRT_STATS
+
+
+# Check environment variable or config for NRT enablement
+ENABLE_NRT = _readenv("NUMBA_CUDA_ENABLE_NRT", bool, False) or getattr(
+    config, "NUMBA_CUDA_ENABLE_NRT", False
+)
+if not hasattr(config, "NUMBA_CUDA_ENABLE_NRT"):
+    config.CUDA_ENABLE_NRT = ENABLE_NRT
+
+
 def get_include():
     """Return the include path for the NRT header"""
     return os.path.dirname(os.path.abspath(__file__))

numba_cuda/numba/cuda/typeconv/rules.py

@@ -4,7 +4,6 @@
 import itertools
 from .typeconv import TypeManager, TypeCastingRules
 from numba.core import types
-from numba.cuda import config
 
 
 default_type_manager = TypeManager()
@@ -16,58 +15,49 @@ def dump_number_rules():
         print(a, "->", b, tm.check_compatible(a, b))
 
 
-if config.USE_LEGACY_TYPE_SYSTEM:  # Old type system
+def _init_casting_rules(tm):
+    tcr = TypeCastingRules(tm)
+    tcr.safe_unsafe(types.boolean, types.int8)
+    tcr.safe_unsafe(types.boolean, types.uint8)
 
-    def _init_casting_rules(tm):
-        tcr = TypeCastingRules(tm)
-        tcr.safe_unsafe(types.boolean, types.int8)
-        tcr.safe_unsafe(types.boolean, types.uint8)
+    tcr.promote_unsafe(types.int8, types.int16)
+    tcr.promote_unsafe(types.uint8, types.uint16)
 
-        tcr.promote_unsafe(types.int8, types.int16)
-        tcr.promote_unsafe(types.uint8, types.uint16)
+    tcr.promote_unsafe(types.int16, types.int32)
+    tcr.promote_unsafe(types.uint16, types.uint32)
 
-        tcr.promote_unsafe(types.int16, types.int32)
-        tcr.promote_unsafe(types.uint16, types.uint32)
+    tcr.promote_unsafe(types.int32, types.int64)
+    tcr.promote_unsafe(types.uint32, types.uint64)
 
-        tcr.promote_unsafe(types.int32, types.int64)
-        tcr.promote_unsafe(types.uint32, types.uint64)
+    tcr.safe_unsafe(types.uint8, types.int16)
+    tcr.safe_unsafe(types.uint16, types.int32)
+    tcr.safe_unsafe(types.uint32, types.int64)
 
-        tcr.safe_unsafe(types.uint8, types.int16)
-        tcr.safe_unsafe(types.uint16, types.int32)
-        tcr.safe_unsafe(types.uint32, types.int64)
+    tcr.safe_unsafe(types.int8, types.float16)
+    tcr.safe_unsafe(types.int16, types.float32)
+    tcr.safe_unsafe(types.int32, types.float64)
 
-        tcr.safe_unsafe(types.int8, types.float16)
-        tcr.safe_unsafe(types.int16, types.float32)
-        tcr.safe_unsafe(types.int32, types.float64)
+    tcr.unsafe_unsafe(types.int16, types.float16)
+    tcr.unsafe_unsafe(types.int32, types.float32)
+    # XXX this is inconsistent with the above; but we want to prefer
+    # float64 over int64 when typing a heterogeneous operation,
+    # e.g. `float64 + int64`.  Perhaps we need more granularity in the
+    # conversion kinds.
+    tcr.safe_unsafe(types.int64, types.float64)
+    tcr.safe_unsafe(types.uint64, types.float64)
 
-        tcr.unsafe_unsafe(types.int16, types.float16)
-        tcr.unsafe_unsafe(types.int32, types.float32)
-        # XXX this is inconsistent with the above; but we want to prefer
-        # float64 over int64 when typing a heterogeneous operation,
-        # e.g. `float64 + int64`.  Perhaps we need more granularity in the
-        # conversion kinds.
-        tcr.safe_unsafe(types.int64, types.float64)
-        tcr.safe_unsafe(types.uint64, types.float64)
+    tcr.promote_unsafe(types.float16, types.float32)
+    tcr.promote_unsafe(types.float32, types.float64)
 
-        tcr.promote_unsafe(types.float16, types.float32)
-        tcr.promote_unsafe(types.float32, types.float64)
+    tcr.safe(types.float32, types.complex64)
+    tcr.safe(types.float64, types.complex128)
 
-        tcr.safe(types.float32, types.complex64)
-        tcr.safe(types.float64, types.complex128)
+    tcr.promote_unsafe(types.complex64, types.complex128)
 
-        tcr.promote_unsafe(types.complex64, types.complex128)
+    # Allow integers to cast ot void*
+    tcr.unsafe_unsafe(types.uintp, types.voidptr)
 
-        # Allow integers to cast ot void*
-        tcr.unsafe_unsafe(types.uintp, types.voidptr)
-
-        return tcr
-else:  # New type system
-    # Currently left as empty
-    # If no casting rules are required we may opt to remove
-    # this framework upon deprecation
-    def _init_casting_rules(tm):
-        tcr = TypeCastingRules(tm)
-        return tcr
+    return tcr
 
 
 default_casting_rules = _init_casting_rules(default_type_manager)