Fix double-registration of items from npydecl Historically cudadecl registered a few things from npydecl to support NumPy dtypes in kernels and a small set of ufuncs. Now that the real npydecl is registered, these items should not also be registered in cudadecl, because this leads to an erroneous double-registration.

Co-authored-by: Graham Markall <[email protected]>

Author: brandon-b-miller

numba_cuda/numba/cuda/cudadecl.py

@@ -6,12 +6,6 @@
 from numba.cuda.typing.npydecl import (
     parse_dtype,
     parse_shape,
-    register_number_classes,
-    register_numpy_ufunc,
-    trigonometric_functions,
-    comparison_functions,
-    math_operations,
-    bit_twiddling_functions,
 )
 from numba.cuda.typing.templates import (
     AttributeTemplate,
@@ -29,8 +23,6 @@
 register_attr = registry.register_attr
 register_global = registry.register_global
 
-register_number_classes(register_global)
-
 
 class Cuda_array_decl(CallableTemplate):
     def generic(self):
@@ -562,19 +554,3 @@ def resolve_local(self, mod):
 
 
 register_global(cuda, types.Module(cuda))
-
-
-# NumPy
-
-for func in trigonometric_functions:
-    register_numpy_ufunc(func, register_global)
-
-for func in comparison_functions:
-    register_numpy_ufunc(func, register_global)
-
-for func in bit_twiddling_functions:
-    register_numpy_ufunc(func, register_global)
-
-for func in math_operations:
-    if func in ("log", "log2", "log10"):
-        register_numpy_ufunc(func, register_global)

numba_cuda/numba/cuda/device_init.py

@@ -87,8 +87,8 @@
     PTXSource,
 )
 
-# from numba.cuda.misc.special import literal_unroll
-# from numba.cuda.misc import literal
+from numba.cuda.misc.special import literal_unroll
+from numba.cuda.misc import literal
 
 reduce = Reduce = reduction.Reduce
 

numba_cuda/numba/cuda/target.py

@@ -36,7 +36,7 @@
 # Typing
 
 
-class CUDATypingContext(typing.Context):
+class CUDATypingContext(typing.BaseContext):
     def load_additional_registries(self):
         from . import (
             cudadecl,
@@ -46,18 +46,18 @@ def load_additional_registries(self):
             libdevicedecl,
             vector_types,
         )
-        from numba.cuda.typing import enumdecl, cffi_utils
+        from numba.cuda.typing import enumdecl, cffi_utils, npydecl
 
         self.install_registry(cudadecl.registry)
         self.install_registry(cffi_utils.registry)
         self.install_registry(cudamath.registry)
         self.install_registry(cmathdecl.registry)
         self.install_registry(libdevicedecl.registry)
+        self.install_registry(npydecl.registry)
         self.install_registry(enumdecl.registry)
         self.install_registry(vector_types.typing_registry)
         self.install_registry(fp16.typing_registry)
         self.install_registry(bf16.typing_registry)
-        super().load_additional_registries()
 
     def resolve_value_type(self, val):
         # treat other dispatcher object as another device function

numba_cuda/numba/cuda/typing/context.py

@@ -522,7 +522,7 @@ def is_external(obj):
             else:
                 # A type was already inserted, see if we can add to it
                 newty = existing.augment(gty)
-                if newty is None and existing != gty:
+                if newty is None:
                     raise TypeError(
                         "cannot augment %s with %s" % (existing, gty)
                     )