Vendor in debuginfo, utils, cffi_utils, enumdecl for CUDA-specific changes

numba_cuda/numba/cuda/cgutils.py

@@ -11,8 +11,8 @@
 
 from llvmlite import ir
 
-from numba.core import types, debuginfo
-from numba.cuda import config, utils
+from numba.core import types
+from numba.cuda import config, utils, debuginfo
 import numba.core.datamodel
 
 

numba_cuda/numba/cuda/debuginfo.py

@@ -1,15 +1,97 @@
 # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 # SPDX-License-Identifier: BSD-2-Clause
 
+import abc
 import os
+from contextlib import contextmanager
 
 from llvmlite import ir
 from numba.core import types, config
 from numba.cuda import cgutils
 from numba.core.datamodel.models import ComplexModel, UnionModel, UniTupleModel
-from numba.core.debuginfo import AbstractDIBuilder
 from numba.cuda.types import GridGroup
 
+
+@contextmanager
+def suspend_emission(builder):
+    """Suspends the emission of debug_metadata for the duration of the context
+    managed block."""
+    ref = builder.debug_metadata
+    builder.debug_metadata = None
+    try:
+        yield
+    finally:
+        builder.debug_metadata = ref
+
+
+class AbstractDIBuilder(metaclass=abc.ABCMeta):
+    @abc.abstractmethod
+    def mark_variable(
+        self,
+        builder,
+        allocavalue,
+        name,
+        lltype,
+        size,
+        line,
+        datamodel=None,
+        argidx=None,
+    ):
+        """Emit debug info for the variable."""
+        pass
+
+    @abc.abstractmethod
+    def mark_location(self, builder, line):
+        """Emit source location information to the given IRBuilder."""
+        pass
+
+    @abc.abstractmethod
+    def mark_subprogram(self, function, qualname, argnames, argtypes, line):
+        """Emit source location information for the given function."""
+        pass
+
+    @abc.abstractmethod
+    def initialize(self):
+        """Initialize the debug info. An opportunity for the debuginfo to
+        prepare any necessary data structures.
+        """
+
+    @abc.abstractmethod
+    def finalize(self):
+        """Finalize the debuginfo by emitting all necessary metadata."""
+        pass
+
+
+class DummyDIBuilder(AbstractDIBuilder):
+    def __init__(self, module, filepath, cgctx, directives_only):
+        pass
+
+    def mark_variable(
+        self,
+        builder,
+        allocavalue,
+        name,
+        lltype,
+        size,
+        line,
+        datamodel=None,
+        argidx=None,
+    ):
+        pass
+
+    def mark_location(self, builder, line):
+        pass
+
+    def mark_subprogram(self, function, qualname, argnames, argtypes, line):
+        pass
+
+    def initialize(self):
+        pass
+
+    def finalize(self):
+        pass
+
+
 _BYTE_SIZE = 8
 
 

numba_cuda/numba/cuda/lowering.py

@@ -10,17 +10,15 @@
 
 from numba.core import (
     typing,
-    utils,
     types,
     ir,
-    debuginfo,
     funcdesc,
     generators,
     config,
     removerefctpass,
     targetconfig,
 )
-from numba.cuda import cgutils
+from numba.cuda import debuginfo, cgutils, utils
 from numba.cuda.core import ir_utils
 from numba.core.errors import (
     LoweringError,

numba_cuda/numba/cuda/target.py

@@ -41,7 +41,7 @@ def load_additional_registries(self):
             libdevicedecl,
             vector_types,
         )
-        from numba.core.typing import enumdecl, cffi_utils
+        from numba.cuda.typing import enumdecl, cffi_utils
 
         self.install_registry(cudadecl.registry)
         self.install_registry(cffi_utils.registry)

numba_cuda/numba/cuda/tests/core/test_serialize.py

@@ -12,7 +12,7 @@
 
 import numba
 from numba.core.errors import TypingError
-from numba.tests.support import TestCase
+from numba.cuda.tests.support import TestCase
 from numba.core.target_extension import resolve_dispatcher_from_str
 from numba.cuda.cloudpickle import dumps, loads
 

numba_cuda/numba/cuda/tests/support.py

@@ -22,7 +22,7 @@
 
 from numba import types
 from numba.core import errors, config
-from numba.core.typing import cffi_utils
+from numba.cuda.typing import cffi_utils
 from numba.cuda.memory_management.nrt import rtsys
 from numba.core.extending import (
     typeof_impl,

numba_cuda/numba/cuda/typing/cffi_utils.py

@@ -0,0 +1,55 @@
+# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: BSD-2-Clause
+
+"""
+Support for CFFI. Allows checking whether objects are CFFI functions and
+obtaining the pointer and numba signature.
+"""
+
+from numba.core import types
+from numba.core.errors import TypingError
+from numba.cuda.typing import templates
+
+try:
+    import cffi
+
+    ffi = cffi.FFI()
+except ImportError:
+    ffi = None
+
+SUPPORTED = ffi is not None
+registry = templates.Registry()
+
+
+@registry.register
+class FFI_from_buffer(templates.AbstractTemplate):
+    key = "ffi.from_buffer"
+
+    def generic(self, args, kws):
+        if kws or len(args) != 1:
+            return
+        [ary] = args
+        if not isinstance(ary, types.Buffer):
+            raise TypingError(
+                "from_buffer() expected a buffer object, got %s" % (ary,)
+            )
+        if ary.layout not in ("C", "F"):
+            raise TypingError(
+                "from_buffer() unsupported on non-contiguous buffers (got %s)"
+                % (ary,)
+            )
+        if ary.layout != "C" and ary.ndim > 1:
+            raise TypingError(
+                "from_buffer() only supports multidimensional arrays with C layout (got %s)"
+                % (ary,)
+            )
+        ptr = types.CPointer(ary.dtype)
+        return templates.signature(ptr, ary)
+
+
+@registry.register_attr
+class FFIAttribute(templates.AttributeTemplate):
+    key = types.ffi
+
+    def resolve_from_buffer(self, ffi):
+        return types.BoundFunction(FFI_from_buffer, types.ffi)

numba_cuda/numba/cuda/typing/enumdecl.py

@@ -0,0 +1,74 @@
+# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: BSD-2-Clause
+
+"""
+Typing for enums.
+"""
+
+import operator
+from numba.core import types
+from numba.cuda.typing.templates import (
+    AbstractTemplate,
+    AttributeTemplate,
+    signature,
+    Registry,
+)
+
+registry = Registry()
+infer = registry.register
+infer_global = registry.register_global
+infer_getattr = registry.register_attr
+
+
+@infer_getattr
+class EnumAttribute(AttributeTemplate):
+    key = types.EnumMember
+
+    def resolve_value(self, ty):
+        return ty.dtype
+
+
+@infer_getattr
+class EnumClassAttribute(AttributeTemplate):
+    key = types.EnumClass
+
+    def generic_resolve(self, ty, attr):
+        """
+        Resolve attributes of an enum class as enum members.
+        """
+        if attr in ty.instance_class.__members__:
+            return ty.member_type
+
+
+@infer
+class EnumClassStaticGetItem(AbstractTemplate):
+    key = "static_getitem"
+
+    def generic(self, args, kws):
+        enum, idx = args
+        if (
+            isinstance(enum, types.EnumClass)
+            and idx in enum.instance_class.__members__
+        ):
+            return signature(enum.member_type, *args)
+
+
+class EnumCompare(AbstractTemplate):
+    def generic(self, args, kws):
+        [lhs, rhs] = args
+        if (
+            isinstance(lhs, types.EnumMember)
+            and isinstance(rhs, types.EnumMember)
+            and lhs == rhs
+        ):
+            return signature(types.boolean, lhs, rhs)
+
+
+@infer_global(operator.eq)
+class EnumEq(EnumCompare):
+    pass
+
+
+@infer_global(operator.ne)
+class EnumNe(EnumCompare):
+    pass

numba_cuda/numba/cuda/typing/templates.py

@@ -15,13 +15,14 @@
 from types import MethodType, FunctionType, MappingProxyType
 
 import numba
-from numba.core import types, utils, targetconfig
+from numba.core import types, targetconfig
 from numba.core.errors import (
     TypingError,
     InternalError,
 )
 from numba.core.cpu_options import InlineOptions
 from numba.core.typing.templates import Signature as CoreSignature
+from numba.cuda import utils
 from numba.cuda.core import ir_utils
 
 # info store for inliner callback functions e.g. cost model

numba_cuda/numba/cuda/utils.py

@@ -29,14 +29,16 @@
 from inspect import Parameter as pyParameter  # noqa: F401
 
 from numba.core.config import (
-    PYVERSION,  # noqa: F401
     MACHINE_BITS,  # noqa: F401
     DEVELOPER_MODE,
 )  # noqa: F401
 from numba.core import types, config
 
 from collections.abc import Mapping, Sequence, MutableSet, MutableMapping
 
+# Python version in (major, minor) tuple
+PYVERSION = sys.version_info[:2]
+
 
 def erase_traceback(exc_value):
     """