Merge branch 'main' into vk/target_extension

VijayKandiah · web-flow · commit f2be4768fb2f · 2025-10-29T13:20:24.000-05:00
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/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/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/cudapy/test_ipc.py b/numba_cuda/numba/cuda/tests/cudapy/test_ipc.py
@@ -23,6 +23,9 @@
 import unittest
 
 
+FUTURE_TIMEOUT = 4
+
+
 def base_ipc_handle_test(handle, size, parent_pid):
     pid = os.getpid()
     assert pid != parent_pid
@@ -97,7 +100,7 @@ def test_ipc_handle(self):
         fut = self.exe.submit(
             base_ipc_handle_test, handle_bytes, size, parent_pid=os.getpid()
         )
-        out = fut.result(timeout=3)
+        out = fut.result(timeout=FUTURE_TIMEOUT)
         np.testing.assert_equal(arr, out)
 
     def variants(self):
@@ -134,7 +137,7 @@ def check_ipc_handle_serialization(self, index_arg=None, foreign=False):
         fut = self.exe.submit(
             serialize_ipc_handle_test, ipch, parent_pid=os.getpid()
         )
-        out = fut.result(timeout=3)
+        out = fut.result(timeout=FUTURE_TIMEOUT)
         np.testing.assert_equal(expect, out)
 
     def test_ipc_handle_serialization(self):
@@ -159,7 +162,7 @@ def check_ipc_array(self, index_arg=None, foreign=False):
 
         # spawn new process for testing
         fut = self.exe.submit(ipc_array_test, ipch, parent_pid=os.getpid())
-        out = fut.result(timeout=3)
+        out = fut.result(timeout=FUTURE_TIMEOUT)
         np.testing.assert_equal(expect, out)
 
     def test_ipc_array(self):
@@ -229,7 +232,9 @@ def test_staged(self):
             for device_num in range(ngpus)
         ]
 
-        for fut in concurrent.futures.as_completed(futures, timeout=3 * ngpus):
+        for fut in concurrent.futures.as_completed(
+            futures, timeout=FUTURE_TIMEOUT * ngpus
+        ):
             np.testing.assert_equal(arr, fut.result())
 
     def test_ipc_array(self):
@@ -243,7 +248,7 @@ def test_ipc_array(self):
             fut = self.exe.submit(
                 staged_ipc_array_test, ipch, device_num, parent_pid=os.getpid()
             )
-            out = fut.result(timeout=3)
+            out = fut.result(timeout=FUTURE_TIMEOUT)
             np.testing.assert_equal(arr, out)
 
 
diff --git a/numba_cuda/numba/cuda/tests/nocuda/test_import.py b/numba_cuda/numba/cuda/tests/nocuda/test_import.py
@@ -38,7 +38,6 @@ def test_no_impl_import(self):
             "numba.cuda.python.printimpl",
             "numba.cuda.core.optional",
             "numba.cuda.misc.gdb_hook",
-            "numba.cuda.misc.literal",
             "numba.cuda.misc.cffiimpl",
             "numba.np.linalg",
             "numba.np.polynomial",
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/ufuncs.py b/numba_cuda/numba/cuda/ufuncs.py
