From 890fcedf4f45006f7f528d84c2f7873a6d0f2d63 Mon Sep 17 00:00:00 2001 From: Vijay Kandiah Date: Tue, 12 Aug 2025 17:02:29 -0700 Subject: [PATCH 1/2] [Refactor] Vendor in debuginfo, utils, cffi_utils, enumdecl for CUDA-specific changes. --- numba_cuda/numba/cuda/cgutils.py | 4 +- numba_cuda/numba/cuda/debuginfo.py | 84 ++++++++++++++++++- numba_cuda/numba/cuda/lowering.py | 4 +- numba_cuda/numba/cuda/target.py | 2 +- .../numba/cuda/tests/core/test_serialize.py | 2 +- numba_cuda/numba/cuda/tests/support.py | 2 +- numba_cuda/numba/cuda/typing/cffi_utils.py | 52 ++++++++++++ numba_cuda/numba/cuda/typing/enumdecl.py | 71 ++++++++++++++++ numba_cuda/numba/cuda/typing/templates.py | 3 +- numba_cuda/numba/cuda/utils.py | 3 + 10 files changed, 217 insertions(+), 10 deletions(-) create mode 100644 numba_cuda/numba/cuda/typing/cffi_utils.py create mode 100644 numba_cuda/numba/cuda/typing/enumdecl.py diff --git a/numba_cuda/numba/cuda/cgutils.py b/numba_cuda/numba/cuda/cgutils.py index 0427fa500..9864a863d 100644 --- a/numba_cuda/numba/cuda/cgutils.py +++ b/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 diff --git a/numba_cuda/numba/cuda/debuginfo.py b/numba_cuda/numba/cuda/debuginfo.py index 54195cd9e..d2fa46b45 100644 --- a/numba_cuda/numba/cuda/debuginfo.py +++ b/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 diff --git a/numba_cuda/numba/cuda/lowering.py b/numba_cuda/numba/cuda/lowering.py index dcf4066ca..a42307a42 100644 --- a/numba_cuda/numba/cuda/lowering.py +++ b/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, diff --git a/numba_cuda/numba/cuda/target.py b/numba_cuda/numba/cuda/target.py index 38077f12f..f321056aa 100644 --- a/numba_cuda/numba/cuda/target.py +++ b/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) diff --git a/numba_cuda/numba/cuda/tests/core/test_serialize.py b/numba_cuda/numba/cuda/tests/core/test_serialize.py index 78bc193fd..65c2e71e1 100644 --- a/numba_cuda/numba/cuda/tests/core/test_serialize.py +++ b/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 diff --git a/numba_cuda/numba/cuda/tests/support.py b/numba_cuda/numba/cuda/tests/support.py index 1c79c0fe3..2e08e1816 100644 --- a/numba_cuda/numba/cuda/tests/support.py +++ b/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, diff --git a/numba_cuda/numba/cuda/typing/cffi_utils.py b/numba_cuda/numba/cuda/typing/cffi_utils.py new file mode 100644 index 000000000..c095c2ee6 --- /dev/null +++ b/numba_cuda/numba/cuda/typing/cffi_utils.py @@ -0,0 +1,52 @@ +""" +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) diff --git a/numba_cuda/numba/cuda/typing/enumdecl.py b/numba_cuda/numba/cuda/typing/enumdecl.py new file mode 100644 index 000000000..e007fb876 --- /dev/null +++ b/numba_cuda/numba/cuda/typing/enumdecl.py @@ -0,0 +1,71 @@ +""" +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 diff --git a/numba_cuda/numba/cuda/typing/templates.py b/numba_cuda/numba/cuda/typing/templates.py index 04f5a3682..43bc3a4cc 100644 --- a/numba_cuda/numba/cuda/typing/templates.py +++ b/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 diff --git a/numba_cuda/numba/cuda/utils.py b/numba_cuda/numba/cuda/utils.py index 05ff5a2d7..ff37a5892 100644 --- a/numba_cuda/numba/cuda/utils.py +++ b/numba_cuda/numba/cuda/utils.py @@ -37,6 +37,9 @@ from collections.abc import Mapping, Sequence, MutableSet, MutableMapping +# Python version in (major, minor) tuple +PYVERSION = sys.version_info[:2] + def erase_traceback(exc_value): """ From 09be92f5e82977ccb1722244dfda87441810df8f Mon Sep 17 00:00:00 2001 From: Vijay Kandiah Date: Thu, 28 Aug 2025 09:49:35 -0700 Subject: [PATCH 2/2] Add SPDX license to new files --- numba_cuda/numba/cuda/typing/cffi_utils.py | 3 +++ numba_cuda/numba/cuda/typing/enumdecl.py | 3 +++ numba_cuda/numba/cuda/utils.py | 1 - 3 files changed, 6 insertions(+), 1 deletion(-) diff --git a/numba_cuda/numba/cuda/typing/cffi_utils.py b/numba_cuda/numba/cuda/typing/cffi_utils.py index c095c2ee6..6f4b78978 100644 --- a/numba_cuda/numba/cuda/typing/cffi_utils.py +++ b/numba_cuda/numba/cuda/typing/cffi_utils.py @@ -1,3 +1,6 @@ +# 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. diff --git a/numba_cuda/numba/cuda/typing/enumdecl.py b/numba_cuda/numba/cuda/typing/enumdecl.py index e007fb876..10a899262 100644 --- a/numba_cuda/numba/cuda/typing/enumdecl.py +++ b/numba_cuda/numba/cuda/typing/enumdecl.py @@ -1,3 +1,6 @@ +# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: BSD-2-Clause + """ Typing for enums. """ diff --git a/numba_cuda/numba/cuda/utils.py b/numba_cuda/numba/cuda/utils.py index ff37a5892..2b1851b98 100644 --- a/numba_cuda/numba/cuda/utils.py +++ b/numba_cuda/numba/cuda/utils.py @@ -29,7 +29,6 @@ from inspect import Parameter as pyParameter # noqa: F401 from numba.core.config import ( - PYVERSION, # noqa: F401 MACHINE_BITS, # noqa: F401 DEVELOPER_MODE, ) # noqa: F401