Vendor in types and datamodel for CUDA-specific changes

numba_cuda/numba/cuda/__init__.py

@@ -7,6 +7,12 @@
 import warnings
 import sys
 
+# Re-export types itself
+import numba.cuda.types as types
+
+# Re-export all type names
+from numba.cuda.types import *
+
 
 # Require NVIDIA CUDA bindings at import time
 if not (

numba_cuda/numba/cuda/_internal/cuda_bf16.py

@@ -18,8 +18,8 @@
 
 import numba
 from llvmlite import ir
-from numba import types
-from numba.core.datamodel import PrimitiveModel, StructModel
+from numba.cuda import types
+from numba.cuda.datamodel import PrimitiveModel, StructModel
 from numba.cuda.extending import (
     lower_cast,
     make_attribute_wrapper,
@@ -41,7 +41,7 @@
 from numba.cuda import CUSource, declare_device
 from numba.cuda.vector_types import vector_types
 from numba.cuda.extending import as_numba_type
-from numba.types import (
+from numba.cuda.types import (
     CPointer,
     Function,
     Number,
@@ -60,7 +60,7 @@
     uint64,
     void,
 )
-from numba.cuda.types import bfloat16
+from numba.cuda.ext_types import bfloat16
 
 float32x2 = vector_types["float32x2"]
 __half = float16

numba_cuda/numba/cuda/_internal/cuda_fp16.py

@@ -18,9 +18,9 @@
 
 import numba
 from llvmlite import ir
-from numba import types
+from numba.cuda import types
 from numba.cuda.cudadrv.driver import _have_nvjitlink
-from numba.core.datamodel import PrimitiveModel, StructModel
+from numba.cuda.datamodel import PrimitiveModel, StructModel
 from numba.core.errors import NumbaPerformanceWarning
 from numba.cuda.extending import (
     lower_cast,
@@ -40,7 +40,7 @@
 from numba.cuda.typing.templates import Registry as TypingRegistry
 from numba.cuda.vector_types import vector_types
 from numba.cuda.extending import as_numba_type
-from numba.types import (
+from numba.cuda.types import (
     CPointer,
     Function,
     Number,
@@ -221,7 +221,7 @@ class _ctor_template_unnamed1362180(ConcreteTemplate):
 
 register_global(unnamed1362180, Function(_ctor_template_unnamed1362180))
 
-__half = _type___half = numba.core.types.float16
+__half = _type___half = numba.cuda.types.float16
 setattr(__half, "alignof_", 2)
 setattr(__half, "align", 2)
 

numba_cuda/numba/cuda/cg.py

@@ -1,12 +1,12 @@
 # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 # SPDX-License-Identifier: BSD-2-Clause
 
-from numba.core import types
+from numba.cuda import types
 from numba.cuda.extending import overload, overload_method
 from numba.cuda.typing import signature
 from numba.cuda import nvvmutils
 from numba.cuda.extending import intrinsic
-from numba.cuda.types import grid_group, GridGroup as GridGroupClass
+from numba.cuda.ext_types import grid_group, GridGroup as GridGroupClass
 
 
 class GridGroup:

numba_cuda/numba/cuda/cgutils.py

@@ -11,9 +11,9 @@
 
 from llvmlite import ir
 
-from numba.core import types
+from numba.cuda import types
 from numba.cuda import config, utils, debuginfo
-import numba.core.datamodel
+import numba.cuda.datamodel
 
 
 bool_t = ir.IntType(1)
@@ -104,7 +104,7 @@ class _StructProxy(object):
     def __init__(self, context, builder, value=None, ref=None):
         self._context = context
         self._datamodel = self._context.data_model_manager[self._fe_type]
-        if not isinstance(self._datamodel, numba.core.datamodel.StructModel):
+        if not isinstance(self._datamodel, numba.cuda.datamodel.StructModel):
             raise TypeError(
                 "Not a structure model: {0}".format(self._datamodel)
             )

numba_cuda/numba/cuda/compiler.py

@@ -7,10 +7,8 @@
 import copy
 
 from numba.core import ir as numba_ir
-from numba.core import (
-    types,
-    bytecode,
-)
+from numba.core import bytecode
+from numba.cuda import types
 from numba.cuda.core.options import ParallelOptions
 from numba.core.compiler_lock import global_compiler_lock
 from numba.core.errors import NumbaWarning, NumbaInvalidConfigWarning

numba_cuda/numba/cuda/core/analysis.py

@@ -2,7 +2,7 @@
 # SPDX-License-Identifier: BSD-2-Clause
 
 from collections import namedtuple, defaultdict
-from numba import types
+from numba.cuda import types
 from numba.core import ir, errors
 from numba.cuda.core import consts
 import operator

numba_cuda/numba/cuda/core/base.py

@@ -11,13 +11,9 @@
 from llvmlite import ir as llvmir
 from llvmlite.ir import Constant
 
-from numba.core import (
-    types,
-    datamodel,
-)
-from numba.cuda import cgutils, debuginfo, utils, config
+from numba.cuda.core import imputils, targetconfig, funcdesc
+from numba.cuda import cgutils, debuginfo, types, utils, datamodel, config
 from numba.core import errors
-from numba.cuda.core import targetconfig, funcdesc, imputils
 from numba.core.compiler_lock import global_compiler_lock
 from numba.cuda.core.pythonapi import PythonAPI
 from numba.cuda.core.imputils import (

numba_cuda/numba/cuda/core/boxing.py

@@ -7,8 +7,7 @@
 
 from llvmlite import ir
 
-from numba.core import types
-from numba.cuda import cgutils
+from numba.cuda import types, cgutils
 from numba.cuda.core.pythonapi import box, unbox, reflect, NativeValue
 from numba.core.errors import NumbaNotImplementedError, TypingError
 from numba.cuda.typing.typeof import typeof, Purpose

numba_cuda/numba/cuda/core/callconv.py

@@ -1,7 +1,7 @@
 # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 # SPDX-License-Identifier: BSD-2-Clause
 
-from numba.core import types
+from numba.cuda import types
 from numba.cuda import cgutils
 from collections import namedtuple
 

numba_cuda/numba/cuda/core/funcdesc.py

@@ -8,7 +8,7 @@
 from collections import defaultdict
 import importlib
 
-from numba.core import types
+from numba.cuda import types
 from numba.cuda import itanium_mangler
 from numba.cuda.utils import _dynamic_modname, _dynamic_module