From 07bd314870ebaf1ac1465e61f59095427cd74c73 Mon Sep 17 00:00:00 2001 From: Vijay Kandiah Date: Wed, 1 Oct 2025 11:40:51 -0700 Subject: [PATCH 1/2] Use numba.core.config directly when available, otherwise use numba.cuda.core.config --- ci/test_thirdparty.sh | 6 +- numba_cuda/numba/cuda/core/config.py | 142 +++--------------- .../numba/cuda/memory_management/nrt.py | 18 ++- 3 files changed, 45 insertions(+), 121 deletions(-) diff --git a/ci/test_thirdparty.sh b/ci/test_thirdparty.sh index 0650f7b5b..cc3ba6f15 100755 --- a/ci/test_thirdparty.sh +++ b/ci/test_thirdparty.sh @@ -39,13 +39,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 diff --git a/numba_cuda/numba/cuda/core/config.py b/numba_cuda/numba/cuda/core/config.py index b7adebba7..ef3b643ae 100644 --- a/numba_cuda/numba/cuda/core/config.py +++ b/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.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 diff --git a/numba_cuda/numba/cuda/memory_management/nrt.py b/numba_cuda/numba/cuda/memory_management/nrt.py index dd0777754..c3a6fc41a 100644 --- a/numba_cuda/numba/cuda/memory_management/nrt.py +++ b/numba_cuda/numba/cuda/memory_management/nrt.py @@ -20,13 +20,29 @@ ) 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 from numba.core.extending import intrinsic, overload_classmethod +# 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__)) From fd125e601907931782dab5b2b655bad59bfba788 Mon Sep 17 00:00:00 2001 From: Vijay Kandiah Date: Fri, 10 Oct 2025 09:53:46 -0700 Subject: [PATCH 2/2] Remove config.USE_LEGACY_TYPE_SYSTEM checks --- numba_cuda/numba/cuda/typeconv/rules.py | 74 +++++++++++-------------- 1 file changed, 32 insertions(+), 42 deletions(-) diff --git a/numba_cuda/numba/cuda/typeconv/rules.py b/numba_cuda/numba/cuda/typeconv/rules.py index 2fa5c1158..8f6513957 100644 --- a/numba_cuda/numba/cuda/typeconv/rules.py +++ b/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)