From 7cb76d1f3722b06841a785ece0140af981481d82 Mon Sep 17 00:00:00 2001 From: Graham Markall Date: Wed, 21 Feb 2024 13:09:42 +0000 Subject: [PATCH 1/7] `BaseContext`: Only attempt to register overloads that match the current target This moves the target check from inside the `_OverloadMethodTemplate` to the `BaseContext.install_registry()` method. The effect of this change is that instead of attempting to register all overloaded method templates with all targets and then silently failing inside the registration, we only attempt to register overloaded methods that are appropriate for the target. The `_OverloadAttributeTemplate` never contained this check, which appears curious at first sight - however, there is a related bug in overloading of attributes that causes their metadata to be lost, and therefore all attributes get registered for all targets regardless of whether they are for the current target or not. This error will be addressed in the following commit. --- numba/core/typing/context.py | 21 +++++++++++++++++++++ numba/core/typing/templates.py | 31 ++++++++++++------------------- 2 files changed, 33 insertions(+), 19 deletions(-) diff --git a/numba/core/typing/context.py b/numba/core/typing/context.py index 5b180c0c606..6c0fadef74b 100644 --- a/numba/core/typing/context.py +++ b/numba/core/typing/context.py @@ -435,9 +435,30 @@ def install_registry(self, registry): except KeyError: loader = templates.RegistryLoader(registry) self._registries[registry] = loader + + from numba.core.target_extension import (get_local_target, + resolve_target_str) + current_target = get_local_target(self) + + def is_for_this_target(ftcls): + metadata = getattr(ftcls, 'metadata', None) + if metadata is None: + return True + + target_str = metadata.get('target') + if target_str is None: + return True + + ft_target = resolve_target_str(target_str) + return current_target.inherits_from(ft_target) + for ftcls in loader.new_registrations('functions'): + if not is_for_this_target(ftcls): + continue self.insert_function(ftcls(self)) for ftcls in loader.new_registrations('attributes'): + if not is_for_this_target(ftcls): + continue self.insert_attributes(ftcls(self)) for gv, gty in loader.new_registrations('globals'): existing = self._lookup_global(gv) diff --git a/numba/core/typing/templates.py b/numba/core/typing/templates.py index 74097fdc0bb..e304599e4f6 100644 --- a/numba/core/typing/templates.py +++ b/numba/core/typing/templates.py @@ -16,7 +16,6 @@ from numba.core.errors import ( TypingError, InternalError, - InternalTargetMismatchError, ) from numba.core.cpu_options import InlineOptions @@ -1096,24 +1095,18 @@ def _init_once(self): """ attr = self._attr - try: - registry = self._get_target_registry('method') - except InternalTargetMismatchError: - # Target mismatch. Do not register attribute lookup here. - pass - else: - lower_builtin = registry.lower - - @lower_builtin((self.key, attr), self.key, types.VarArg(types.Any)) - def method_impl(context, builder, sig, args): - typ = sig.args[0] - typing_context = context.typing_context - fnty = self._get_function_type(typing_context, typ) - sig = self._get_signature(typing_context, fnty, sig.args, {}) - call = context.get_function(fnty, sig) - # Link dependent library - context.add_linking_libs(getattr(call, 'libs', ())) - return call(builder, args) + registry = self._get_target_registry('method') + + @registry.lower((self.key, attr), self.key, types.VarArg(types.Any)) + def method_impl(context, builder, sig, args): + typ = sig.args[0] + typing_context = context.typing_context + fnty = self._get_function_type(typing_context, typ) + sig = self._get_signature(typing_context, fnty, sig.args, {}) + call = context.get_function(fnty, sig) + # Link dependent library + context.add_linking_libs(getattr(call, 'libs', ())) + return call(builder, args) def _resolve(self, typ, attr): if self._attr != attr: From 7cee0ea3e6c47adf01b1d367124863f7650247c2 Mon Sep 17 00:00:00 2001 From: Graham Markall Date: Wed, 21 Feb 2024 16:23:32 +0000 Subject: [PATCH 2/7] Fix loss of metadata for overload_attribute Decorator keywords arguments other than `inline` were dropped, so other metadata such as the target were silently lost. We resolve this by passing all keyword arguments to `make_overload_attribute_template()`. Also add a test that verifies that the metadata still exists, which prevents a CUDA-only attribute overload being used on the CPU target, with a sensible TypingError as the result. --- numba/core/extending.py | 2 +- numba/core/typing/templates.py | 2 +- numba/cuda/tests/cudapy/test_overload.py | 30 +++++++++++++++++++++--- 3 files changed, 29 insertions(+), 5 deletions(-) diff --git a/numba/core/extending.py b/numba/core/extending.py index 48339598d0e..e4da9109988 100644 --- a/numba/core/extending.py +++ b/numba/core/extending.py @@ -190,7 +190,7 @@ def get(arr): def decorate(overload_func): template = make_overload_attribute_template( typ, attr, overload_func, - inline=kwargs.get('inline', 'never'), + **kwargs ) infer_getattr(template) overload(overload_func, **kwargs)(overload_func) diff --git a/numba/core/typing/templates.py b/numba/core/typing/templates.py index e304599e4f6..2b371cfe720 100644 --- a/numba/core/typing/templates.py +++ b/numba/core/typing/templates.py @@ -1155,7 +1155,7 @@ def get_template_info(self): return types.BoundFunction(MethodTemplate, typ) -def make_overload_attribute_template(typ, attr, overload_func, inline, +def make_overload_attribute_template(typ, attr, overload_func, inline='never', prefer_literal=False, base=_OverloadAttributeTemplate, **kwargs): diff --git a/numba/cuda/tests/cudapy/test_overload.py b/numba/cuda/tests/cudapy/test_overload.py index 2764678ec6a..4c19201a052 100644 --- a/numba/cuda/tests/cudapy/test_overload.py +++ b/numba/cuda/tests/cudapy/test_overload.py @@ -1,7 +1,8 @@ -from numba import cuda, njit -from numba.core.extending import overload +from numba import cuda, njit, types +from numba.core.errors import TypingError +from numba.core.extending import overload, overload_attribute from numba.cuda.testing import CUDATestCase, skip_on_cudasim, unittest - +from numba.tests.test_extending import mydummy_type, MyDummyType import numpy as np @@ -295,6 +296,29 @@ def kernel(x): expected = GENERIC_TARGET_OL_CALLS_TARGET_OL * GENERIC_TARGET_OL self.check_overload_cpu(kernel, expected) + def test_overload_attribute_target(self): + @overload_attribute(MyDummyType, 'cuda_only', target='cuda') + def ov_dummy_cuda_attr(obj): + def imp(obj): + return 42 + + return imp + + # Ensure that we cannot use the CUDA target-specific attribute on the + # CPU, and that an appropriate typing error is raised + with self.assertRaisesRegex(TypingError, + "Unknown attribute 'cuda_only'"): + @njit(types.void(mydummy_type)) + def illegal_target_attr_use(x): + return x.cuda_only + + # Ensure that the CUDA target-specific attribute is usable and works + # correctly when the target is CUDA - note eager compilation via + # signature + @cuda.jit(types.void(types.int64[::1], mydummy_type)) + def cuda_target_attr_use(res, dummy): + res[0] = dummy.cuda_only + if __name__ == '__main__': unittest.main() From bbbf8f5d0d9ebae2ac323a44639cad101848e523 Mon Sep 17 00:00:00 2001 From: Graham Markall Date: Wed, 10 Apr 2024 10:09:16 +0100 Subject: [PATCH 3/7] Add a target extension test ensuring metadata is present and correct This is analagous to the similar test added for CUDA in the previous commit, but doesn't require CUDA to run so can be a part of the main test suite. --- numba/tests/test_target_extension.py | 27 ++++++++++++++++++++++++++- 1 file changed, 26 insertions(+), 1 deletion(-) diff --git a/numba/tests/test_target_extension.py b/numba/tests/test_target_extension.py index c472b1c71c7..3157819da72 100644 --- a/numba/tests/test_target_extension.py +++ b/numba/tests/test_target_extension.py @@ -12,7 +12,8 @@ from functools import cached_property import numpy as np from numba import njit, types -from numba.extending import overload, intrinsic, overload_classmethod +from numba.extending import (overload, overload_attribute, + overload_classmethod, intrinsic) from numba.core.target_extension import ( JitDecorator, target_registry, @@ -41,6 +42,7 @@ from numba.core.compiler import CompilerBase, DefaultPassBuilder from numba.core.compiler_machinery import FunctionPass, register_pass from numba.core.typed_passes import PreLowerStripPhis +from numba.tests.test_extending import mydummy_type, MyDummyType # Define a new target, this target extends GPU, this places the DPU in the # target hierarchy as a type of GPU. @@ -704,6 +706,29 @@ def foo(): from numba.core.runtime import nrt self.assertIsInstance(r, nrt.MemInfo) + def test_overload_attribute_target(self): + @overload_attribute(MyDummyType, 'dpu_only', target='dpu') + def ov_dummy_dpu_attr(obj): + def imp(obj): + return 42 + + return imp + + # Ensure that we cannot use the DPU target-specific attribute on the + # CPU, and that an appropriate typing error is raised + with self.assertRaisesRegex(errors.TypingError, + "Unknown attribute 'dpu_only'"): + @njit(types.void(mydummy_type)) + def illegal_target_attr_use(x): + return x.dpu_only + + # Ensure that the DPU target-specific attribute is usable and works + # correctly when the target is DPU - note eager compilation via + # signature + @djit(types.void(types.int64[::1], mydummy_type)) + def cuda_target_attr_use(res, dummy): + res[0] = dummy.dpu_only + class TestTargetOffload(TestCase): """In this use case the CPU compilation pipeline is extended with a new From bd1732c3689143ec0b55207fc05a4ac9302c9667 Mon Sep 17 00:00:00 2001 From: Graham Markall Date: Thu, 11 Apr 2024 11:04:52 +0100 Subject: [PATCH 4/7] Target extension tests: fix DPU target setting The DPU target should have been setting a target override during compilation (similar to how the CUDA target does), and it should also set the `target_backend` target option, mirroring what the default `@jit` decorator does. Not setting these doesn't seem to have had any ill effect within the testsuite, but these settings should both be made to avoid running into latent issues in future. --- numba/tests/test_target_extension.py | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/numba/tests/test_target_extension.py b/numba/tests/test_target_extension.py index 3157819da72..d0f626d4715 100644 --- a/numba/tests/test_target_extension.py +++ b/numba/tests/test_target_extension.py @@ -283,6 +283,10 @@ def typing_context(self): class DPUDispatcher(Dispatcher): targetdescr = dpu_target + def compile(self, sig): + with target_override('dpu'): + return super().compile(sig) + # Register a dispatcher for the DPU target, a lot of the code uses this # internally to work out what to do RE compilation @@ -319,6 +323,8 @@ def dispatcher_wrapper(self): if "nopython" in self._kwargs: topt["nopython"] = True + topt['target_backend'] = 'dpu' + # It would be easy to specialise the default compilation pipeline for # this target here. pipeline_class = compiler.Compiler From a6fd2d0486e18e9aa1883727bff9e932a498d42c Mon Sep 17 00:00:00 2001 From: Graham Markall Date: Thu, 11 Apr 2024 11:07:38 +0100 Subject: [PATCH 5/7] Ignore registrations with invalid target during typing context refresh When typing contexts are refreshed, there may be registrations for invalid targets pending (e.g. from other tests that exercise error handling in the presence of invalid user code). These registrations need to be ignored - see the comments in the code for further detail. Attempting to compile a call to an overload with an invalid target now causes a `NonexistentTargetError` to be raised, so the target extension tests are updated accordingly. --- numba/core/errors.py | 6 ++++++ numba/core/target_extension.py | 5 +++-- numba/core/typing/context.py | 25 ++++++++++++++++++++++++- numba/tests/test_target_extension.py | 4 ++-- 4 files changed, 35 insertions(+), 5 deletions(-) diff --git a/numba/core/errors.py b/numba/core/errors.py index 0d339583ff8..892b230bd33 100644 --- a/numba/core/errors.py +++ b/numba/core/errors.py @@ -720,6 +720,12 @@ def __init__(self, kind, target_hw, hw_clazz): super().__init__(msg) +class NonexistentTargetError(InternalError): + """For signalling that a target that does not exist was requested. + """ + pass + + class RequireLiteralValue(TypingError): """ For signalling that a function's typing requires a constant value for diff --git a/numba/core/target_extension.py b/numba/core/target_extension.py index d0546671c64..91e5df09985 100644 --- a/numba/core/target_extension.py +++ b/numba/core/target_extension.py @@ -1,7 +1,8 @@ from abc import ABC, abstractmethod from numba.core.registry import DelayedRegistry, CPUDispatcher from numba.core.decorators import jit -from numba.core.errors import InternalTargetMismatchError, NumbaValueError +from numba.core.errors import (InternalTargetMismatchError, + NonexistentTargetError) from threading import local as tls @@ -18,7 +19,7 @@ def __getitem__(self, item): msg = "No target is registered against '{}', known targets:\n{}" known = '\n'.join([f"{k: <{10}} -> {v}" for k, v in target_registry.items()]) - raise NumbaValueError(msg.format(item, known)) from None + raise NonexistentTargetError(msg.format(item, known)) from None # Registry mapping target name strings to Target classes diff --git a/numba/core/typing/context.py b/numba/core/typing/context.py index 6c0fadef74b..dd50eca923d 100644 --- a/numba/core/typing/context.py +++ b/numba/core/typing/context.py @@ -449,7 +449,30 @@ def is_for_this_target(ftcls): if target_str is None: return True - ft_target = resolve_target_str(target_str) + # There may be pending registrations for nonexistent targets. + # Ideally it would be impossible to leave a registration pending + # for an invalid target, but in practice this is exceedingly + # difficult to guard against - many things are registered at import + # time, and eagerly reporting an error when registering for invalid + # targets would require that all target registration code is + # executed prior to all typing registrations during the import + # process; attempting to enforce this would impose constraints on + # execution order during import that would be very difficult to + # resolve and maintain in the presence of typical code maintenance. + # Furthermore, these constraints would be imposed not only on + # Numba internals, but also on its dependents. + # + # Instead of that enforcement, we simply catch any occurrences of + # registrations for targets that don't exist, and report that + # they're not for this target. They will then not be encountered + # again during future typing context refreshes (because the + # loader's new registrations are a stream_list that doesn't yield + # previously-yielded items). + try: + ft_target = resolve_target_str(target_str) + except errors.NonexistentTargetError: + return False + return current_target.inherits_from(ft_target) for ftcls in loader.new_registrations('functions'): diff --git a/numba/tests/test_target_extension.py b/numba/tests/test_target_extension.py index d0f626d4715..212a138a035 100644 --- a/numba/tests/test_target_extension.py +++ b/numba/tests/test_target_extension.py @@ -520,7 +520,7 @@ def dpu_foo(): def test_invalid_target_jit(self): - with self.assertRaises(errors.NumbaValueError) as raises: + with self.assertRaises(errors.NonexistentTargetError) as raises: @njit(_target='invalid_silicon') def foo(): pass @@ -537,7 +537,7 @@ def bar(): # This is a typing error at present as it fails during typing when the # overloads are walked. - with self.assertRaises(errors.TypingError) as raises: + with self.assertRaises(errors.NonexistentTargetError) as raises: @overload(bar, target='invalid_silicon') def ol_bar(): return lambda : None From 8a8cefe4a2545079bb3b87843c43d1b564c8e189 Mon Sep 17 00:00:00 2001 From: Graham Markall Date: Wed, 24 Apr 2024 22:46:32 +0100 Subject: [PATCH 6/7] Remove target_backend option for DPU target This option is going to be removed by #9539 when it is merged. --- numba/tests/test_target_extension.py | 2 -- 1 file changed, 2 deletions(-) diff --git a/numba/tests/test_target_extension.py b/numba/tests/test_target_extension.py index 212a138a035..f3fea74e660 100644 --- a/numba/tests/test_target_extension.py +++ b/numba/tests/test_target_extension.py @@ -323,8 +323,6 @@ def dispatcher_wrapper(self): if "nopython" in self._kwargs: topt["nopython"] = True - topt['target_backend'] = 'dpu' - # It would be easy to specialise the default compilation pipeline for # this target here. pipeline_class = compiler.Compiler From 0d823b3989e64cc3297e6de3813c24b75815a42e Mon Sep 17 00:00:00 2001 From: Graham Markall Date: Wed, 24 Apr 2024 23:00:38 +0100 Subject: [PATCH 7/7] Fix test cross-contamination in attribute overload tests Instead of reusing `MyDummyType` and `mydummy_type` from `test_extending`, we use the `make_dummy_type` method of the test classes to create types that are unique to the current test. A minor (inconsequential) error in the signatures in the tests is also fixed. --- numba/cuda/tests/cudapy/test_overload.py | 7 +++++-- numba/tests/test_target_extension.py | 7 +++++-- 2 files changed, 10 insertions(+), 4 deletions(-) diff --git a/numba/cuda/tests/cudapy/test_overload.py b/numba/cuda/tests/cudapy/test_overload.py index 4c19201a052..412fe2434ca 100644 --- a/numba/cuda/tests/cudapy/test_overload.py +++ b/numba/cuda/tests/cudapy/test_overload.py @@ -1,8 +1,8 @@ from numba import cuda, njit, types from numba.core.errors import TypingError from numba.core.extending import overload, overload_attribute +from numba.core.typing.typeof import typeof from numba.cuda.testing import CUDATestCase, skip_on_cudasim, unittest -from numba.tests.test_extending import mydummy_type, MyDummyType import numpy as np @@ -297,6 +297,9 @@ def kernel(x): self.check_overload_cpu(kernel, expected) def test_overload_attribute_target(self): + MyDummy, MyDummyType = self.make_dummy_type() + mydummy_type = typeof(MyDummy()) + @overload_attribute(MyDummyType, 'cuda_only', target='cuda') def ov_dummy_cuda_attr(obj): def imp(obj): @@ -308,7 +311,7 @@ def imp(obj): # CPU, and that an appropriate typing error is raised with self.assertRaisesRegex(TypingError, "Unknown attribute 'cuda_only'"): - @njit(types.void(mydummy_type)) + @njit(types.int64(mydummy_type)) def illegal_target_attr_use(x): return x.cuda_only diff --git a/numba/tests/test_target_extension.py b/numba/tests/test_target_extension.py index f3fea74e660..a78b234b81e 100644 --- a/numba/tests/test_target_extension.py +++ b/numba/tests/test_target_extension.py @@ -33,6 +33,7 @@ from numba.core.codegen import CPUCodegen, JITCodeLibrary from numba.core.callwrapper import PyCallWrapper from numba.core.imputils import RegistryLoader, Registry +from numba.core.typing.typeof import typeof from numba import _dynfunc import llvmlite.binding as ll from llvmlite import ir as llir @@ -42,7 +43,6 @@ from numba.core.compiler import CompilerBase, DefaultPassBuilder from numba.core.compiler_machinery import FunctionPass, register_pass from numba.core.typed_passes import PreLowerStripPhis -from numba.tests.test_extending import mydummy_type, MyDummyType # Define a new target, this target extends GPU, this places the DPU in the # target hierarchy as a type of GPU. @@ -711,6 +711,9 @@ def foo(): self.assertIsInstance(r, nrt.MemInfo) def test_overload_attribute_target(self): + MyDummy, MyDummyType = self.make_dummy_type() + mydummy_type = typeof(MyDummy()) + @overload_attribute(MyDummyType, 'dpu_only', target='dpu') def ov_dummy_dpu_attr(obj): def imp(obj): @@ -722,7 +725,7 @@ def imp(obj): # CPU, and that an appropriate typing error is raised with self.assertRaisesRegex(errors.TypingError, "Unknown attribute 'dpu_only'"): - @njit(types.void(mydummy_type)) + @njit(types.int64(mydummy_type)) def illegal_target_attr_use(x): return x.dpu_only