Eagerly use all vended numba.cuda modules, Bump supported Numba-CUDA version to 0.21+

ci/test_conda_python.sh

@@ -23,7 +23,7 @@ rapids-mamba-retry create -n test \
   cuda-version=${RAPIDS_CUDA_VERSION%.*} \
   cuda-nvrtc \
   numba >=0.59 \
-  "numba-cuda>=0.20.1,<0.21.0" \
+  "numba-cuda>=0.21.0,<0.23.0" \
   cuda-cudart-dev \
   python=${RAPIDS_PY_VERSION} \
   cffi \

conda/environment-cu12.yaml

@@ -9,7 +9,7 @@ dependencies:
   - python>=3.10
   - cmake>=3.29 # Allow overriding CMAKE_PREFIX_PATH and CMAKE_INSTALL_PREFIX
   - clangdev >=18,<22.0
-  - numba-cuda[cu12] >=0.20.1,<0.21.0
+  - numba-cuda[cu12] >=0.21.0,<0.23.0
   - pybind11
   - pytest
   - pytest-benchmark

conda/environment-cu13.yaml

@@ -9,7 +9,7 @@ dependencies:
   - python>=3.10
   - cmake>=3.29 # Allow overriding CMAKE_PREFIX_PATH and CMAKE_INSTALL_PREFIX
   - clangdev >=18,<22.0
-  - numba-cuda[cu13] >=0.20.1,<0.21.0
+  - numba-cuda[cu13] >=0.21.0,<0.23.0
   - pybind11
   - pytest
   - pytest-benchmark

conda/environment_template.yaml

@@ -9,7 +9,7 @@ dependencies:
   - python={{ python_version }}
   - cmake>=3.29 # Allow overriding CMAKE_PREFIX_PATH and CMAKE_INSTALL_PREFIX
   - clangdev >=18,<22.0
-  - numba-cuda >=0.20.1,<0.21.0
+  - numba-cuda >=0.21.0,<0.23.0
   - pybind11
   - pytest
   - pytest-benchmark

conda/recipes/numbast/meta.yaml

@@ -52,7 +52,7 @@ requirements:
     - cuda-version >=12.5
     - numba >=0.59
     - python
-    - numba-cuda >=0.20.1,<0.21.0
+    - numba-cuda >=0.21.0,<0.23.0
     - pyyaml
     - click
     - jinja2

conda/recipes/numbast_extensions/meta.yaml

@@ -43,7 +43,7 @@ requirements:
     - cuda-cudart-dev
     - python
     - numba >=0.59
-    - numba-cuda >=0.20.1,<0.21.0
+    - numba-cuda >=0.21.0,<0.23.0
     - numbast >=0.2.0
 
   test:

docs/source/faq.rst

@@ -122,3 +122,48 @@ Notes
   libc++, ensure its headers are installed (e.g., ``/usr/include/c++/v1`` or Conda ``libcxx-devel``) and pass
   ``-stdlib=libc++``; otherwise libstdc++ is typically selected by default on Linux.
 - AST Canopy is linked against Clang 20. For host resources, please use corresponding Clang version.
+
+
+Generated bindings and Numba-CUDA version requirements
+-------------------------------------------------------
+
+**What version of numba-cuda do generated bindings require?**
+
+Bindings generated by Numbast have specific version requirements for ``numba-cuda`` at runtime. The version of
+Numbast used to generate the bindings determines the compatible ``numba-cuda`` versions.
+
+.. list-table:: Numbast to numba-cuda compatibility
+   :header-rows: 1
+   :widths: 30 70
+
+   * - Numbast Version
+     - Required numba-cuda Version
+   * - 0.6.0 (current dev)
+     - ``>=0.21.0,<0.23.0``
+   * - 0.5.x
+     - ``>=0.20.1,<0.21.0``
+
+**Why do generated bindings have version requirements?**
+
+Numbast generates Python code that uses Numba-CUDA's internal APIs. These APIs can change between releases, so
+bindings generated with a specific version of Numbast are tested against a specific range of ``numba-cuda`` versions.
+
+**How do I ensure compatibility?**
+
+*For dynamic binding generation:*
+
+- The correct ``numba-cuda`` version constraints are automatically enforced at the package dependency level
+  and managed by your package manager (pip or conda). When you install Numbast, compatible versions of
+  ``numba-cuda`` are installed automatically via the dependencies specified in ``pyproject.toml`` and Conda
+  environment files.
+
+*For static binding generation:*
+
+- When distributing generated bindings, document the required ``numba-cuda`` version range in your package
+  dependencies so users can install a compatible version.
+- Generated static bindings (see :doc:`Static binding generation <static>`) can be regenerated with newer
+  versions of Numbast if you need to support newer ``numba-cuda`` releases.
+
+.. note::
+   These version restrictions may be relaxed or removed once ``numba-cuda`` releases a stable 1.0 version with
+   stabilized public APIs. Until then, bindings are tested against specific version ranges to ensure compatibility.

numbast/benchmarks/test_arithmetic.py

@@ -1,6 +1,6 @@
 import numba.cuda as cuda
+from numba.cuda.types import float32
 import numpy as np
-from numba import float32
 
 import pytest
 

numbast/pyproject.toml

@@ -18,7 +18,7 @@ classifiers = [
 ]
 dependencies = [
   "numba>=0.59.0",
-  "numba-cuda>=0.20.1,<0.21.0",
+  "numba-cuda>=0.21.0,<0.23.0",
   "ast_canopy>=0.5.0",
   "pyyaml",
   "click",
@@ -41,12 +41,12 @@ repository = "https://github.com/NVIDIA/numbast"
 [project.optional-dependencies]
 dev = ["ruff"]
 docs = ["sphinx>=7.0", "sphinx-copybutton>=0.5.2"]
-test-cu12 = ["pytest", "cffi", "numba-cuda[cu12]>=0.20.1,<0.21.0"]
+test-cu12 = ["pytest", "cffi", "numba-cuda[cu12]>=0.21.0,<0.23.0"]
 test-cu13 = [
   "pytest",
   "cffi",
   "cuda-toolkit[cudart,crt,curand,cccl,nvcc]==13.*",
-  "numba-cuda[cu13]>=0.20.1,<0.21.0",
+  "numba-cuda[cu13]>=0.21.0,<0.23.0",
 ]
 
 

numbast/src/numbast/class_template.py

@@ -24,7 +24,7 @@
 from numba.cuda import declare_device
 from numba.cuda.cudadecl import register_global, register, register_attr
 from numba.cuda.cudaimpl import lower
-from numba.core.imputils import numba_typeref_ctor
+from numba.cuda.core.imputils import numba_typeref_ctor
 from numba.core.typing.npydecl import parse_dtype
 from numba.core.errors import RequireLiteralValue, TypingError
 

numbast/src/numbast/function.py

@@ -5,7 +5,7 @@
 from collections import defaultdict
 
 from numba import types as nbtypes
-from numba.core.typing import signature as nb_signature, Signature
+from numba.cuda.typing import signature as nb_signature, Signature
 from numba.cuda.typing.templates import ConcreteTemplate
 from numba.cuda import declare_device
 from numba.cuda.cudadecl import register_global, register

numbast/src/numbast/functor.py

@@ -2,8 +2,8 @@
 # SPDX-License-Identifier: Apache-2.0
 
 from numba import types
-from numba.core.extending import typeof_impl
-from numba.core.imputils import lower_constant
+from numba.cuda.extending import typeof_impl
+from numba.cuda.core.imputils import lower_constant
 from numba import cuda
 
 from numbast.types import NUMBA_TO_CTYPE_MAPS as N2C, FunctorType

numbast/src/numbast/static/function.py

@@ -227,11 +227,11 @@ def _render_decl_device(self):
         """Render codes that declares a foreign function for this function in Numba."""
 
         self.Imports.add("from numba.cuda import declare_device")
-        self.Imports.add("from numba.core.typing import signature")
+        self.Imports.add("from numba.cuda.typing import signature")
         # All arguments are passed by pointers in C-CPP shim interop
-        self.Imports.add("from numba.types import CPointer")
+        self.Imports.add("from numba.cuda.types import CPointer")
         # Numba ABI returns int32 for exception codes
-        self.Imports.add("from numba.types import int32")
+        self.Imports.add("from numba.cuda.types import int32")
 
         decl_device_rendered = self.decl_device_template.format(
             decl_name=self._deduplicated_shim_name,
@@ -273,7 +273,7 @@ def _render_shim_function(self):
     def _render_lowering(self):
         """Render lowering codes for this struct constructor."""
 
-        self.Imports.add("from numba.core.typing import signature")
+        self.Imports.add("from numba.cuda.typing import signature")
 
         use_cooperative = ""
         if self._use_cooperative:

numbast/src/numbast/static/renderer.py

@@ -122,8 +122,8 @@ def _try_import_numba_type(cls, typ: str):
             cls.Imported_VectorTypes.append(typ)
             cls._imported_numba_types.add(typ)
 
-        elif typ in numba.types.__dict__:
-            cls.Imports.add(f"from numba.types import {typ}")
+        elif typ in numba.cuda.types.__dict__:
+            cls.Imports.add(f"from numba.cuda.types import {typ}")
             cls._imported_numba_types.add(typ)
 
         else:
@@ -348,7 +348,7 @@ def registry_setup(use_separate_registry: bool) -> str:
             "from numba.cuda.typing.templates import Registry as TypingRegistry"
         )
         BaseRenderer.Imports.add(
-            "from numba.core.imputils import Registry as TargetRegistry"
+            "from numba.cuda.core.imputils import Registry as TargetRegistry"
         )
         return BaseRenderer.SeparateRegistrySetup
     else:
@@ -364,5 +364,5 @@ def registry_setup(use_separate_registry: bool) -> str:
         BaseRenderer.Imports.add(
             "from numba.cuda.cudaimpl import lower_constant"
         )
-        BaseRenderer.Imports.add("from numba.core.extending import lower_cast")
+        BaseRenderer.Imports.add("from numba.cuda.extending import lower_cast")
         return ""

numbast/src/numbast/static/struct.py

@@ -7,8 +7,8 @@
 import tempfile
 import warnings
 
-from numba.types import Type
-from numba.core.datamodel.models import StructModel, PrimitiveModel
+from numba.cuda.types import Type
+from numba.cuda.datamodel.models import StructModel, PrimitiveModel
 
 from ast_canopy.pylibastcanopy import access_kind, method_kind
 from ast_canopy.decl import Struct, StructMethod
@@ -216,11 +216,11 @@ def _render_decl_device(self):
         """Render codes that declares a foreign function for this constructor in Numba."""
 
         self.Imports.add("from numba.cuda import declare_device")
-        self.Imports.add("from numba.core.typing import signature")
+        self.Imports.add("from numba.cuda.typing import signature")
         # All arguments are passed by pointers in C-CPP shim interop
-        self.Imports.add("from numba.types import CPointer")
+        self.Imports.add("from numba.cuda.types import CPointer")
         # Numba ABI returns int32 for exception codes
-        self.Imports.add("from numba.types import int32")
+        self.Imports.add("from numba.cuda.types import int32")
 
         decl_device_rendered = self.struct_ctor_decl_device_template.format(
             struct_ctor_device_decl_str=self._struct_ctor_device_decl_str,
@@ -396,7 +396,7 @@ def _render_typing(self, signature_strs: list[str]):
         self.Imports.add(
             "from numba.cuda.typing.templates import ConcreteTemplate"
         )
-        self.Imports.add("from numba.types import Function")
+        self.Imports.add("from numba.cuda.types import Function")
 
         signatures_str = ", ".join(signature_strs)
 
@@ -565,9 +565,9 @@ def _render_decl_device(self):
         """Render codes that declares a foreign function for this constructor in Numba."""
 
         self.Imports.add("from numba.cuda import declare_device")
-        self.Imports.add("from numba.core.typing import signature")
+        self.Imports.add("from numba.cuda.typing import signature")
         # All arguments are passed by pointers in C-CPP shim interop
-        self.Imports.add("from numba.types import CPointer")
+        self.Imports.add("from numba.cuda.types import CPointer")
 
         decl_device_rendered = (
             self.struct_conversion_op_decl_device_template.format(
@@ -1101,7 +1101,7 @@ def __init__(self):
         self.bitwidth = {struct_sizeof} * 8
 
     def can_convert_from(self, typingctx, other):
-        from numba.core.typeconv import Conversion
+        from numba.cuda.typeconv import Conversion
         if other in [{implicit_conversion_types}]:
             return Conversion.safe
 
@@ -1204,12 +1204,12 @@ def __init__(
         self._data_model = data_model
 
         self.Imports.add(
-            f"from numba.types import {self._parent_type.__qualname__}"
+            f"from numba.cuda.types import {self._parent_type.__qualname__}"
         )
         self._parent_type_str = self._parent_type.__qualname__
 
         self.Imports.add(
-            f"from numba.core.datamodel import {self._data_model.__qualname__}"
+            f"from numba.cuda.datamodel import {self._data_model.__qualname__}"
         )
         self._data_model_str = self._data_model.__qualname__
 
@@ -1262,7 +1262,7 @@ def _render_python_api(self):
 
         This is the python handle to use it in Numba kernels.
         """
-        self.Imports.add("from numba.extending import as_numba_type")
+        self.Imports.add("from numba.cuda.extending import as_numba_type")
 
         self._python_api_rendered = self.python_api_template.format(
             struct_type_name=self._struct_type_name,
@@ -1283,7 +1283,7 @@ def _render_data_model(self):
         - Sets self._data_model_rendered.
         """
 
-        self.Imports.add("from numba.core.extending import register_model")
+        self.Imports.add("from numba.cuda.extending import register_model")
 
         if self._data_model == PrimitiveModel:
             self.Imports.add("from llvmlite import ir")
@@ -1325,7 +1325,7 @@ def _render_struct_attr(self):
                 "from numba.cuda.typing.templates import AttributeTemplate"
             )
             self.Imports.add(
-                "from numba.core.extending import make_attribute_wrapper"
+                "from numba.cuda.extending import make_attribute_wrapper"
             )
             # For method attribute resolution
             self.Imports.add("from numba.types import BoundFunction")

numbast/src/numbast/static/tests/test_conflict_shim_names.py

@@ -3,8 +3,8 @@
 
 
 from numba import cuda
-from numba.types import Type
-from numba.core.datamodel import StructModel
+from numba.cuda.types import Type
+from numba.cuda.datamodel import StructModel
 
 
 def _check_shim_name_contains_mangled_name(src: str, mangled_name: str):

numbast/src/numbast/static/tests/test_function_static_bindings.py

@@ -6,7 +6,7 @@
 import numpy as np
 import cffi
 
-from numba.types import int32, float32
+from numba.cuda.types import int32, float32
 from numba import cuda
 from numba.cuda import device_array
 

numbast/src/numbast/static/tests/test_link_two_files.py

@@ -4,8 +4,8 @@
 import pytest
 
 from numba import cuda
-from numba.types import Type, Number
-from numba.core.datamodel import StructModel, PrimitiveModel
+from numba.cuda.types import Type, Number
+from numba.cuda.datamodel import StructModel, PrimitiveModel
 
 from ast_canopy import parse_declarations_from_source
 from numbast.static.renderer import clear_base_renderer_cache, registry_setup

numbast/src/numbast/static/tests/test_operator_bindings.py

@@ -4,8 +4,8 @@
 import pytest
 
 from numba import cuda
-from numba.types import Type
-from numba.core.datamodel import StructModel
+from numba.cuda.types import Type
+from numba.cuda.datamodel import StructModel
 from numba.cuda import device_array
 
 from ast_canopy import parse_declarations_from_source

numbast/src/numbast/static/tests/test_static_demo.py

@@ -3,9 +3,11 @@
 
 import pytest
 
+import numpy as np
+
 from numba import cuda
-from numba.types import Number, float64
-from numba.core.datamodel import PrimitiveModel
+from numba.cuda.types import Number, float64
+from numba.cuda.datamodel import PrimitiveModel
 from numba.cuda import device_array
 
 from ast_canopy import parse_declarations_from_source
@@ -53,4 +55,5 @@ def kernel(arr):
 
     arr = device_array((1,), "float64")
     kernel[1, 1](arr)
-    assert all(arr.copy_to_host() == [])
+    host = arr.copy_to_host()
+    assert np.allclose(host, [3.14], rtol=1e-3, atol=1e-3)

numbast/src/numbast/static/tests/test_struct_static_bindings.py

@@ -6,8 +6,8 @@
 import numpy as np
 
 from numba import cuda
-from numba.types import Type, Number
-from numba.core.datamodel import StructModel, PrimitiveModel
+from numba.cuda.types import Type, Number
+from numba.cuda.datamodel import StructModel, PrimitiveModel
 from numba.cuda import device_array
 
 from ast_canopy import parse_declarations_from_source