Merge branch 'main' into atmn/vendor-in-errors

gmarkall · web-flow · commit 1ba09ab00852 · 2025-10-30T08:07:07.000Z
diff --git a/.github/workflows/conda-python-build.yaml b/.github/workflows/conda-python-build.yaml
@@ -55,6 +55,7 @@ jobs:
     runs-on: ubuntu-latest
     outputs:
       MATRIX: ${{ steps.compute-matrix.outputs.MATRIX }}
+    timeout-minutes: 60
     steps:
       - name: Compute Build Matrix
         id: compute-matrix
@@ -82,6 +83,7 @@ jobs:
       image: rapidsai/ci-conda:cuda${{ matrix.CUDA_VER }}-${{ matrix.LINUX_VER }}-py${{ matrix.PY_VER }}
       env:
         RAPIDS_BUILD_TYPE: ${{ inputs.build_type }}
+    timeout-minutes: 60
     steps:
       - uses: actions/checkout@v4
         with:
diff --git a/.github/workflows/conda-python-tests.yaml b/.github/workflows/conda-python-tests.yaml
@@ -58,6 +58,7 @@ jobs:
       BUILD_TYPE: ${{ inputs.build_type }}
     outputs:
       MATRIX: ${{ steps.compute-matrix.outputs.MATRIX }}
+    timeout-minutes: 60
     steps:
       - name: Compute Python Test Matrix
         id: compute-matrix
@@ -90,6 +91,7 @@ jobs:
       env:
         RAPIDS_BUILD_TYPE: ${{ inputs.build_type }}
         NVIDIA_VISIBLE_DEVICES: ${{ env.NVIDIA_VISIBLE_DEVICES }}
+    timeout-minutes: 60
     steps:
       - uses: actions/checkout@v4
         with:
diff --git a/.github/workflows/docs-build.yaml b/.github/workflows/docs-build.yaml
@@ -52,6 +52,7 @@ jobs:
       image: rapidsai/ci-conda:cuda13.0.0-ubuntu24.04-py3.13
       env:
         RAPIDS_BUILD_TYPE: ${{ inputs.build_type }}
+    timeout-minutes: 60
     steps:
       - uses: actions/checkout@v4
         with:
diff --git a/.github/workflows/simulator-test.yaml b/.github/workflows/simulator-test.yaml
@@ -52,6 +52,7 @@ jobs:
     runs-on: ubuntu-latest
     outputs:
       MATRIX: ${{ steps.compute-matrix.outputs.MATRIX }}
+    timeout-minutes: 60
     steps:
       - name: Compute Simulator Test Matrix
         id: compute-matrix
@@ -79,6 +80,7 @@ jobs:
       image: rapidsai/ci-conda:cuda${{ matrix.CUDA_VER }}-${{ matrix.LINUX_VER }}-py${{ matrix.PY_VER }}
       env:
         RAPIDS_BUILD_TYPE: ${{ inputs.build_type }}
+    timeout-minutes: 60
     steps:
       - uses: actions/checkout@v4
         with:
diff --git a/.github/workflows/wheel-windows-build.yaml b/.github/workflows/wheel-windows-build.yaml
@@ -26,6 +26,7 @@ jobs:
           - "3.12"
           - "3.13"
     runs-on: windows-2022
+    timeout-minutes: 60
     steps:
       - uses: actions/checkout@v4
         with:
diff --git a/.github/workflows/wheel-windows-tests.yaml b/.github/workflows/wheel-windows-tests.yaml
@@ -26,6 +26,7 @@ jobs:
           - "3.12"
           - "3.13"
     runs-on: "cuda-python-windows-gpu-github"
+    timeout-minutes: 60
     steps:
       - uses: actions/checkout@v4
         with:
diff --git a/.github/workflows/wheels-build.yaml b/.github/workflows/wheels-build.yaml
@@ -76,6 +76,7 @@ jobs:
     runs-on: ubuntu-latest
     outputs:
       MATRIX: ${{ steps.compute-matrix.outputs.MATRIX }}
+    timeout-minutes: 60
     steps:
       - name: Compute Build Matrix
         id: compute-matrix
@@ -102,7 +103,7 @@ jobs:
       image: "rapidsai/ci-wheel:cuda${{ matrix.CUDA_VER }}-${{ matrix.LINUX_VER }}-py${{ matrix.PY_VER }}"
       env:
         RAPIDS_BUILD_TYPE: ${{ inputs.build_type }}
-
+    timeout-minutes: 60
     steps:
       - name: checkout code repo
         uses: actions/checkout@v4
diff --git a/.github/workflows/wheels-test.yaml b/.github/workflows/wheels-test.yaml
@@ -67,6 +67,7 @@ jobs:
       BUILD_TYPE: ${{ inputs.build_type }}
     outputs:
       MATRIX: ${{ steps.compute-matrix.outputs.MATRIX }}
+    timeout-minutes: 60
     steps:
       - name: Validate test type
         run: |
@@ -104,6 +105,7 @@ jobs:
       env:
         NVIDIA_VISIBLE_DEVICES: ${{ env.NVIDIA_VISIBLE_DEVICES }} # GPU jobs must set this container env variable
         RAPIDS_BUILD_TYPE: ${{ inputs.build_type }}
+    timeout-minutes: 60
     steps:
     - name: Run nvidia-smi to make sure GPU is working
       run: nvidia-smi
diff --git a/ci/test_wheel.sh b/ci/test_wheel.sh
@@ -13,7 +13,7 @@ echo "Package path: ${package}"
 DEPENDENCIES=(
     "${package}"
     "cuda-python==${CUDA_VER_MAJOR_MINOR%.*}.*"
-    "cuda-core==0.3.*"
+    "cuda-core>=0.3.0,<1.0.0"
     "--group"
     "test"
 )
diff --git a/conda/recipes/numba-cuda/meta.yaml b/conda/recipes/numba-cuda/meta.yaml
@@ -30,7 +30,7 @@ requirements:
     - python
     - numba >=0.59.1
     - cuda-bindings >=12.9.1
-    - cuda-core ==0.3.*
+    - cuda-core >=0.3.0,<1.0.0
 
 about:
   home: {{ project_urls["Homepage"] }}
diff --git a/numba_cuda/numba/cuda/compiler.py b/numba_cuda/numba/cuda/compiler.py
@@ -741,10 +741,7 @@ def compile_cuda(
     flags.max_registers = max_registers
     flags.lto = lto
 
-    # Run compilation pipeline
-    from numba.core.target_extension import target_override
-
-    with target_override("cuda"):
+    with utils.numba_target_override():
         cres = compile_extra(
             typingctx=typingctx,
             targetctx=targetctx,
diff --git a/numba_cuda/numba/cuda/core/base.py b/numba_cuda/numba/cuda/core/base.py
@@ -3,6 +3,7 @@
 
 from collections import defaultdict
 import copy
+import importlib
 import sys
 from itertools import permutations, takewhile
 from contextlib import contextmanager
@@ -212,10 +213,15 @@ def enable_boundscheck(self, value):
     def __init__(self, typing_context, target):
         self.address_size = utils.MACHINE_BITS
         self.typing_context = typing_context
-        from numba.core.target_extension import target_registry
-
         self.target_name = target
-        self.target = target_registry[target]
+
+        if importlib.util.find_spec("numba"):
+            from numba.core.target_extension import CUDA
+
+            # Used only in Numba's target_extension implementation.
+            # Numba-CUDA has the target_extension implementation removed, and
+            # references to it hardcoded to values specific to the CUDA target.
+            self.target = CUDA
 
         # A mapping of installed registries to their loaders
         self._registries = {}
diff --git a/numba_cuda/numba/cuda/cpython/listobj.py b/numba_cuda/numba/cuda/cpython/listobj.py
@@ -300,8 +300,12 @@ def define_dtor(self):
         mod = builder.module
         # Declare dtor
         fnty = ir.FunctionType(ir.VoidType(), [cgutils.voidptr_t])
+        if isinstance(self.dtype, types.containers.List):
+            dtypestr = f"list_{self.dtype.dtype}"
+        else:
+            dtypestr = str(self.dtype)
         fn = cgutils.get_or_insert_function(
-            mod, fnty, ".dtor.list.{}".format(self.dtype)
+            mod, fnty, f"numba_cuda_dtor_list_{dtypestr}"
         )
         if not fn.is_declaration:
             # End early if the dtor is already defined
diff --git a/numba_cuda/numba/cuda/cudadecl.py b/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)
diff --git a/numba_cuda/numba/cuda/device_init.py b/numba_cuda/numba/cuda/device_init.py
@@ -87,6 +87,9 @@
     PTXSource,
 )
 
+from numba.cuda.misc.special import literal_unroll
+from numba.cuda.misc import literal
+
 reduce = Reduce = reduction.Reduce
 
 # Expose vector type constructors and aliases as module level attributes.
diff --git a/numba_cuda/numba/cuda/dispatcher.py b/numba_cuda/numba/cuda/dispatcher.py
@@ -726,12 +726,8 @@ class CUDACache(Cache):
     _impl_class = CUDACacheImpl
 
     def load_overload(self, sig, target_context):
-        # Loading an overload refreshes the context to ensure it is
-        # initialized. To initialize the correct (i.e. CUDA) target, we need to
-        # enforce that the current target is the CUDA target.
-        from numba.core.target_extension import target_override
-
-        with target_override("cuda"):
+        # Loading an overload refreshes the context to ensure it is initialized.
+        with utils.numba_target_override():
             return super().load_overload(sig, target_context)
 
 
diff --git a/numba_cuda/numba/cuda/initialize.py b/numba_cuda/numba/cuda/initialize.py
@@ -5,15 +5,3 @@
 def initialize_all():
     # Import models to register them with the data model manager
     import numba.cuda.models  # noqa: F401
-
-    from numba.cuda.decorators import jit
-    from numba.cuda.dispatcher import CUDADispatcher
-    from numba.core.target_extension import (
-        target_registry,
-        dispatcher_registry,
-        jit_registry,
-    )
-
-    cuda_target = target_registry["cuda"]
-    jit_registry[cuda_target] = jit
-    dispatcher_registry[cuda_target] = CUDADispatcher
diff --git a/numba_cuda/numba/cuda/lowering.py b/numba_cuda/numba/cuda/lowering.py
@@ -1240,12 +1240,9 @@ def _lower_call_normal(self, fnty, expr, signature):
             )
         tname = expr.target
         if tname is not None:
-            from numba.core.target_extension import (
-                resolve_dispatcher_from_str,
-            )
+            from numba.cuda.descriptor import cuda_target
 
-            disp = resolve_dispatcher_from_str(tname)
-            hw_ctx = disp.targetdescr.target_context
+            hw_ctx = cuda_target.target_context
             impl = hw_ctx.get_function(fnty, signature)
         else:
             impl = self.context.get_function(fnty, signature)
diff --git a/numba_cuda/numba/cuda/memory_management/nrt.cu b/numba_cuda/numba/cuda/memory_management/nrt.cu
@@ -78,6 +78,39 @@ extern "C" __device__ void NRT_MemInfo_call_dtor(NRT_MemInfo* mi)
   NRT_MemInfo_destroy(mi);
 }
 
+static void __device__
+nrt_varsize_dtor(void *ptr, size_t size, void *info) {
+    if (info) {
+        /* call element dtor */
+        typedef void dtor_fn_t(void *ptr);
+        dtor_fn_t *dtor = (dtor_fn_t *)info;
+        dtor(ptr);
+    }
+    NRT_Free(ptr);
+}
+
+extern "C"
+__device__ NRT_MemInfo* NRT_MemInfo_new_varsize(size_t size)
+{
+    void *data = NRT_Allocate(size);
+    if (data == NULL) {
+        return NULL; /* return early as allocation failed */
+    }
+
+    NRT_MemInfo *mi = NRT_MemInfo_new(data, size, nrt_varsize_dtor, NULL);
+
+    return mi;
+}
+
+extern "C"
+__device__ NRT_MemInfo* NRT_MemInfo_new_varsize_dtor(size_t size, NRT_dtor_function dtor) {
+    NRT_MemInfo *mi = NRT_MemInfo_new_varsize(size);
+    if (mi) {
+        mi->dtor_info = (void*)dtor;
+    }
+    return mi;
+}
+
 extern "C" __device__ void* NRT_MemInfo_data_fast(NRT_MemInfo *mi)
 {
   return mi->data;
@@ -167,11 +200,13 @@ extern "C" __device__ void NRT_decref(NRT_MemInfo* mi)
 }
 
 
-#endif
+
 
 extern "C" __device__ void NRT_incref(NRT_MemInfo* mi)
 {
   if (mi != NULL) {
     mi->refct++;
   }
 }
+
+#endif
diff --git a/numba_cuda/numba/cuda/memory_management/nrt.cuh b/numba_cuda/numba/cuda/memory_management/nrt.cuh
@@ -44,3 +44,5 @@ extern "C" __device__ void NRT_MemInfo_init(NRT_MemInfo* mi,
                                             size_t size,
                                             NRT_dtor_function dtor,
                                             void* dtor_info);
+extern "C" __device__ NRT_MemInfo* NRT_MemInfo_new_varsize(size_t size);
+extern "C" __device__ NRT_MemInfo* NRT_MemInfo_new_varsize_dtor(size_t size, NRT_dtor_function dtor);
diff --git a/numba_cuda/numba/cuda/target.py b/numba_cuda/numba/cuda/target.py
@@ -46,13 +46,14 @@ 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)
@@ -179,11 +180,8 @@ def load_additional_registries(self):
         from numba.cuda.cpython import builtins as cpython_builtins
         from numba.cuda.core import optional  # noqa: F401
         from numba.cuda.misc import cffiimpl
-        from numba.cuda.np import (
-            arrayobj,
-            npdatetime,
-            polynomial,
-        )
+        from numba.cuda.np import arrayobj, npdatetime, polynomial, arraymath
+
         from . import (
             cudaimpl,
             fp16,
@@ -223,6 +221,7 @@ def load_additional_registries(self):
         self.install_registry(polynomial.registry)
         self.install_registry(npdatetime.registry)
         self.install_registry(arrayobj.registry)
+        self.install_registry(arraymath.registry)
 
         # Install only implementations that are defined outside of numba (i.e.,
         # in third-party extensions) from Numba's builtin_registry.
@@ -240,6 +239,14 @@ def target_data(self):
             self._target_data = ll.create_target_data(nvvm.NVVM().data_layout)
         return self._target_data
 
+    def build_list(self, builder, list_type, items):
+        """
+        Build a list from the Numba *list_type* and its initial *items*.
+        """
+        from numba.cuda.cpython import listobj
+
+        return listobj.build_list(self, builder, list_type, items)
+
     @cached_property
     def nonconst_module_attrs(self):
         """
diff --git a/numba_cuda/numba/cuda/tests/core/test_serialize.py b/numba_cuda/numba/cuda/tests/core/test_serialize.py
diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_ipc.py b/numba_cuda/numba/cuda/tests/cudapy/test_ipc.py
diff --git a/numba_cuda/numba/cuda/tests/nocuda/test_import.py b/numba_cuda/numba/cuda/tests/nocuda/test_import.py
diff --git a/numba_cuda/numba/cuda/tests/test_array_reductions.py b/numba_cuda/numba/cuda/tests/test_array_reductions.py
diff --git a/numba_cuda/numba/cuda/types/cuda_functions.py b/numba_cuda/numba/cuda/types/cuda_functions.py
diff --git a/numba_cuda/numba/cuda/typing/context.py b/numba_cuda/numba/cuda/typing/context.py
diff --git a/numba_cuda/numba/cuda/typing/templates.py b/numba_cuda/numba/cuda/typing/templates.py
diff --git a/numba_cuda/numba/cuda/ufuncs.py b/numba_cuda/numba/cuda/ufuncs.py
diff --git a/numba_cuda/numba/cuda/utils.py b/numba_cuda/numba/cuda/utils.py
diff --git a/pyproject.toml b/pyproject.toml

Original file line number	Diff line number	Diff line change
`@@ -13,7 +13,7 @@ echo "Package path: ${package}"`
`13`	`13`	`DEPENDENCIES=(`
`14`	`14`	`"${package}"`
`15`	`15`	`"cuda-python==${CUDA_VER_MAJOR_MINOR%.}."`
`16`		`- "cuda-core==0.3.*"`
	`16`	`+ "cuda-core>=0.3.0,<1.0.0"`
`17`	`17`	`"--group"`
`18`	`18`	`"test"`
`19`	`19`	`)`