diff --git a/ci/test_conda_python.sh b/ci/test_conda_python.sh index 73c88e17..75a9c7bf 100755 --- a/ci/test_conda_python.sh +++ b/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 \ diff --git a/conda/environment-cu12.yaml b/conda/environment-cu12.yaml index 970b4c12..f8debd67 100644 --- a/conda/environment-cu12.yaml +++ b/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 diff --git a/conda/environment-cu13.yaml b/conda/environment-cu13.yaml index d7be7861..a02c8935 100644 --- a/conda/environment-cu13.yaml +++ b/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 diff --git a/conda/environment_template.yaml b/conda/environment_template.yaml index cbc83926..f9374889 100644 --- a/conda/environment_template.yaml +++ b/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 diff --git a/conda/recipes/numbast/meta.yaml b/conda/recipes/numbast/meta.yaml index 688cf62d..d9794ff5 100644 --- a/conda/recipes/numbast/meta.yaml +++ b/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 diff --git a/conda/recipes/numbast_extensions/meta.yaml b/conda/recipes/numbast_extensions/meta.yaml index a3fcd0ba..b15e924b 100644 --- a/conda/recipes/numbast_extensions/meta.yaml +++ b/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: diff --git a/docs/source/faq.rst b/docs/source/faq.rst index 96a0fc8a..8c694ed7 100644 --- a/docs/source/faq.rst +++ b/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 `) 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. diff --git a/numbast/benchmarks/test_arithmetic.py b/numbast/benchmarks/test_arithmetic.py index e41d8d98..06d7e6dc 100644 --- a/numbast/benchmarks/test_arithmetic.py +++ b/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 diff --git a/numbast/pyproject.toml b/numbast/pyproject.toml index d8394d27..af490037 100644 --- a/numbast/pyproject.toml +++ b/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", ] diff --git a/numbast/src/numbast/class_template.py b/numbast/src/numbast/class_template.py index 60cf81a0..3a37c37e 100644 --- a/numbast/src/numbast/class_template.py +++ b/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 diff --git a/numbast/src/numbast/function.py b/numbast/src/numbast/function.py index b1b4df3d..5e3d3a83 100644 --- a/numbast/src/numbast/function.py +++ b/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 diff --git a/numbast/src/numbast/functor.py b/numbast/src/numbast/functor.py index 552a3d13..c2808cf9 100644 --- a/numbast/src/numbast/functor.py +++ b/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 diff --git a/numbast/src/numbast/static/function.py b/numbast/src/numbast/static/function.py index 20da4ac0..0cd34e6e 100644 --- a/numbast/src/numbast/static/function.py +++ b/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: diff --git a/numbast/src/numbast/static/renderer.py b/numbast/src/numbast/static/renderer.py index 51a5200c..c2842069 100644 --- a/numbast/src/numbast/static/renderer.py +++ b/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 "" diff --git a/numbast/src/numbast/static/struct.py b/numbast/src/numbast/static/struct.py index 41176669..c3a3484d 100644 --- a/numbast/src/numbast/static/struct.py +++ b/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") diff --git a/numbast/src/numbast/static/tests/test_conflict_shim_names.py b/numbast/src/numbast/static/tests/test_conflict_shim_names.py index b621556d..91e7dce1 100644 --- a/numbast/src/numbast/static/tests/test_conflict_shim_names.py +++ b/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): diff --git a/numbast/src/numbast/static/tests/test_function_static_bindings.py b/numbast/src/numbast/static/tests/test_function_static_bindings.py index 750e920f..be1ad953 100644 --- a/numbast/src/numbast/static/tests/test_function_static_bindings.py +++ b/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 diff --git a/numbast/src/numbast/static/tests/test_link_two_files.py b/numbast/src/numbast/static/tests/test_link_two_files.py index 13bc5183..317145b9 100644 --- a/numbast/src/numbast/static/tests/test_link_two_files.py +++ b/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 diff --git a/numbast/src/numbast/static/tests/test_operator_bindings.py b/numbast/src/numbast/static/tests/test_operator_bindings.py index ff22a087..27eddc46 100644 --- a/numbast/src/numbast/static/tests/test_operator_bindings.py +++ b/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 diff --git a/numbast/src/numbast/static/tests/test_static_demo.py b/numbast/src/numbast/static/tests/test_static_demo.py index 058a1a21..210af25a 100644 --- a/numbast/src/numbast/static/tests/test_static_demo.py +++ b/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) diff --git a/numbast/src/numbast/static/tests/test_struct_static_bindings.py b/numbast/src/numbast/static/tests/test_struct_static_bindings.py index 9575e020..a56e5488 100644 --- a/numbast/src/numbast/static/tests/test_struct_static_bindings.py +++ b/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 diff --git a/numbast/src/numbast/struct.py b/numbast/src/numbast/struct.py index f3b0805e..4b4c3a14 100644 --- a/numbast/src/numbast/struct.py +++ b/numbast/src/numbast/struct.py @@ -6,15 +6,15 @@ from llvmlite import ir -from numba import types as nbtypes -from numba.core.extending import ( +from numba.cuda import types as nbtypes +from numba.cuda.extending import ( register_model, lower_cast, make_attribute_wrapper, ) -from numba.core.typing import signature as nb_signature +from numba.cuda.typing import signature as nb_signature from numba.cuda.typing.templates import ConcreteTemplate, AttributeTemplate -from numba.core.datamodel.models import StructModel, PrimitiveModel +from numba.cuda.datamodel.models import StructModel, PrimitiveModel from numba.cuda import declare_device from numba.cuda.cudadecl import register_global, register, register_attr from numba.cuda.cudaimpl import lower diff --git a/numbast/src/numbast/types.py b/numbast/src/numbast/types.py index 43112229..15fcb896 100644 --- a/numbast/src/numbast/types.py +++ b/numbast/src/numbast/types.py @@ -6,8 +6,8 @@ from numba import types as nbtypes from numba.cuda.types import bfloat16 +from numba.cuda.typing.typeof import typeof from numba.cuda.vector_types import vector_types -from numba.misc.special import typeof from numba.cuda._internal.cuda_bf16 import _type_unnamed1405307 diff --git a/numbast/tests/demo/demo.py b/numbast/tests/demo/demo.py index 5a1e017c..5bdb03de 100644 --- a/numbast/tests/demo/demo.py +++ b/numbast/tests/demo/demo.py @@ -3,7 +3,7 @@ from numbast import bind_cxx_struct, bind_cxx_function, MemoryShimWriter from numba import types, cuda -from numba.core.datamodel.models import PrimitiveModel +from numba.cuda.datamodel.models import PrimitiveModel import numpy as np diff --git a/numbast/tests/test_struct.py b/numbast/tests/test_struct.py index 48665daa..0cced722 100644 --- a/numbast/tests/test_struct.py +++ b/numbast/tests/test_struct.py @@ -6,7 +6,7 @@ import numpy as np from numba import types, cuda, float32 -from numba.core.datamodel import StructModel +from numba.cuda.datamodel import StructModel from llvmlite import ir diff --git a/numbast_extensions/pyproject.toml b/numbast_extensions/pyproject.toml index bb8fc9a7..6065c596 100644 --- a/numbast_extensions/pyproject.toml +++ b/numbast_extensions/pyproject.toml @@ -15,7 +15,7 @@ classifiers = [ "License :: OSI Approved :: Apache Software License", "Operating System :: OS Independent", ] -dependencies = ["numba-cuda>=0.20.1,<0.21.0", "numbast>=0.2.0"] +dependencies = ["numba-cuda>=0.21.0,<0.23.0", "numbast>=0.2.0"] [project.urls] homepage = "https://github.com/NVIDIA/numbast"