From 96ffc6743b14a2740b370e51118c8233b93dda53 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 6 Nov 2025 13:45:53 -0800 Subject: [PATCH 1/8] eagerly use all vended numba.cuda modules --- numbast/benchmarks/test_arithmetic.py | 2 +- numbast/src/numbast/function.py | 2 +- numbast/src/numbast/functor.py | 4 +-- numbast/src/numbast/static/function.py | 8 ++--- numbast/src/numbast/static/renderer.py | 6 ++-- numbast/src/numbast/static/struct.py | 30 +++++++++---------- .../static/tests/test_conflict_shim_names.py | 4 +-- .../tests/test_function_static_bindings.py | 2 +- .../static/tests/test_link_two_files.py | 4 +-- .../static/tests/test_operator_bindings.py | 4 +-- .../numbast/static/tests/test_static_demo.py | 4 +-- .../tests/test_struct_static_bindings.py | 4 +-- numbast/src/numbast/struct.py | 8 ++--- numbast/src/numbast/types.py | 2 +- numbast/tests/data/sample_function.cuh | 8 +++++ numbast/tests/demo/demo.py | 2 +- numbast/tests/test_struct.py | 2 +- 17 files changed, 52 insertions(+), 44 deletions(-) create mode 100644 numbast/tests/data/sample_function.cuh 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/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 6c8ef057..35605f06 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 1fbcbc4b..1be50e89 100644 --- a/numbast/src/numbast/static/renderer.py +++ b/numbast/src/numbast/static/renderer.py @@ -104,7 +104,7 @@ def _try_import_numba_type(cls, typ: str): cls._imported_numba_types.add(typ) elif typ in numba.types.__dict__: - cls.Imports.add(f"from numba.types import {typ}") + cls.Imports.add(f"from numba.cuda.types import {typ}") cls._imported_numba_types.add(typ) else: @@ -283,7 +283,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, lower_cast" + "from numba.cuda.core.imputils import Registry as TargetRegistry, lower_cast" ) return BaseRenderer.SeparateRegistrySetup else: @@ -299,5 +299,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 18d70edf..4a7ae9ba 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 @@ -201,11 +201,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, @@ -363,7 +363,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) @@ -522,9 +522,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( @@ -565,7 +565,7 @@ def _render_shim_function(self): def _render_lowering(self): """Render lowering codes for this struct constructor.""" - self.Imports.add("from numba.core.extending import lower_cast") + self.Imports.add("from numba.cuda.extending import lower_cast") self._lowering_rendered = ( self.struct_conversion_op_lowering_template.format( @@ -689,7 +689,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 @@ -764,12 +764,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__ @@ -816,7 +816,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, @@ -826,7 +826,7 @@ def _render_python_api(self): def _render_data_model(self): """Renders the data model of the struct.""" - 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") @@ -868,7 +868,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" ) public_fields = [ 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..99250b57 100644 --- a/numbast/src/numbast/static/tests/test_static_demo.py +++ b/numbast/src/numbast/static/tests/test_static_demo.py @@ -4,8 +4,8 @@ import pytest 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 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 ea93988a..a538b98a 100644 --- a/numbast/src/numbast/static/tests/test_struct_static_bindings.py +++ b/numbast/src/numbast/static/tests/test_struct_static_bindings.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 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 c61dc182..899d83be 100644 --- a/numbast/src/numbast/struct.py +++ b/numbast/src/numbast/struct.py @@ -5,15 +5,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 c0f16faa..f75a26fc 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 class FunctorType(nbtypes.Type): diff --git a/numbast/tests/data/sample_function.cuh b/numbast/tests/data/sample_function.cuh new file mode 100644 index 00000000..9083cc4c --- /dev/null +++ b/numbast/tests/data/sample_function.cuh @@ -0,0 +1,8 @@ +// clang-format off +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// clang-format on + +#pragma once + +void __device__ func_with_void_ptr_arg(void *ptr) {} 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 66ede4a5..f84baf26 100644 --- a/numbast/tests/test_struct.py +++ b/numbast/tests/test_struct.py @@ -4,7 +4,7 @@ import os from numba import types -from numba.core.datamodel import StructModel +from numba.cuda.datamodel import StructModel from llvmlite import ir From df501fe53d1fffc63a457a5f874cdc2a8bd85d38 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 4 Dec 2025 13:14:26 -0800 Subject: [PATCH 2/8] pin numba-cuda versions to >=0.21.0 --- ci/test_conda_python.sh | 2 +- conda/environment-cu12.yaml | 2 +- conda/environment-cu13.yaml | 2 +- conda/environment_template.yaml | 2 +- conda/recipes/numbast/meta.yaml | 2 +- conda/recipes/numbast_extensions/meta.yaml | 2 +- numbast/pyproject.toml | 6 +++--- numbast_extensions/pyproject.toml | 2 +- 8 files changed, 10 insertions(+), 10 deletions(-) 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/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_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" From 5b10afb70816af4532dbff30d44d3ecbbff5284d Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 4 Dec 2025 13:35:55 -0800 Subject: [PATCH 3/8] document version requirements for bindings v.s. numba-cuda --- docs/source/faq.rst | 45 +++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 45 insertions(+) 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. From 0fe0bf048c266e37c3114e10946e28c9c84b4b5e Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 4 Dec 2025 14:53:02 -0800 Subject: [PATCH 4/8] check on numba.cuda.types --- numbast/src/numbast/static/renderer.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/numbast/src/numbast/static/renderer.py b/numbast/src/numbast/static/renderer.py index 1cec1e84..c2842069 100644 --- a/numbast/src/numbast/static/renderer.py +++ b/numbast/src/numbast/static/renderer.py @@ -122,7 +122,7 @@ 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__: + elif typ in numba.cuda.types.__dict__: cls.Imports.add(f"from numba.cuda.types import {typ}") cls._imported_numba_types.add(typ) From fa5f7afa2cde5175ee5b5e8cf46e07f01ff0d0e8 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 4 Dec 2025 15:04:02 -0800 Subject: [PATCH 5/8] avoid unused parameter --- numbast/tests/data/sample_function.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/numbast/tests/data/sample_function.cuh b/numbast/tests/data/sample_function.cuh index 9083cc4c..cc0ba3fc 100644 --- a/numbast/tests/data/sample_function.cuh +++ b/numbast/tests/data/sample_function.cuh @@ -5,4 +5,4 @@ #pragma once -void __device__ func_with_void_ptr_arg(void *ptr) {} +__device__ void func_with_void_ptr_arg(void *ptr) { (void)ptr; } From 0375b487344f1b293b5df15aa0d99211bcdd2496 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 4 Dec 2025 15:25:21 -0800 Subject: [PATCH 6/8] use numba_typeref_ctor from vended module --- numbast/src/numbast/class_template.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 From 56f62d600f5e2b2bd36acd4e2a40588c4df1ff7f Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 4 Dec 2025 16:17:36 -0800 Subject: [PATCH 7/8] fix: bug in static demo test --- numbast/src/numbast/static/tests/test_static_demo.py | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/numbast/src/numbast/static/tests/test_static_demo.py b/numbast/src/numbast/static/tests/test_static_demo.py index 99250b57..210af25a 100644 --- a/numbast/src/numbast/static/tests/test_static_demo.py +++ b/numbast/src/numbast/static/tests/test_static_demo.py @@ -3,6 +3,8 @@ import pytest +import numpy as np + from numba import cuda from numba.cuda.types import Number, float64 from numba.cuda.datamodel import PrimitiveModel @@ -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) From 6b6a52836e4d0ef2a7c5afc953c05d7db42e51c1 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 11 Dec 2025 07:59:15 -0800 Subject: [PATCH 8/8] remove unrelated test file --- numbast/tests/data/sample_function.cuh | 8 -------- 1 file changed, 8 deletions(-) delete mode 100644 numbast/tests/data/sample_function.cuh diff --git a/numbast/tests/data/sample_function.cuh b/numbast/tests/data/sample_function.cuh deleted file mode 100644 index cc0ba3fc..00000000 --- a/numbast/tests/data/sample_function.cuh +++ /dev/null @@ -1,8 +0,0 @@ -// clang-format off -// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -// SPDX-License-Identifier: Apache-2.0 -// clang-format on - -#pragma once - -__device__ void func_with_void_ptr_arg(void *ptr) { (void)ptr; }