Enable global lists within kernels

numba_cuda/numba/cuda/np/arrayobj.py

@@ -3643,6 +3643,8 @@ def constant_array(context, builder, ty, pyval):
     """
     Create a constant array (mechanism is target-dependent).
     """
+    if isinstance(pyval, list):
+        pyval = np.asarray(pyval)
     return context.make_constant_array(builder, ty, pyval)
 
 

numba_cuda/numba/cuda/target.py

@@ -15,6 +15,7 @@
 from numba.cuda.core.base import BaseContext
 from numba.cuda.typing import cmathdecl
 from numba.cuda import datamodel
+from numba.cuda.cpython import listobj
 
 from .cudadrv import nvvm
 from numba.cuda import (
@@ -164,7 +165,6 @@ def load_additional_registries(self):
             numbers,
             slicing,
             iterators,
-            listobj,
             unicode,
             charseq,
             cmathimpl,
@@ -239,7 +239,6 @@ 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)
 

numba_cuda/numba/cuda/tests/cudapy/test_lists.py

@@ -0,0 +1,38 @@
+# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: BSD-2-Clause
+
+"""
+Test cases for a list object within CUDA kernels
+"""
+
+from numba import cuda
+from numba.cuda import config
+import numpy as np
+from numba.cuda.testing import (
+    CUDATestCase,
+)
+
+
+class ListTest(CUDATestCase):
+    def setUp(self):
+        self.old_nrt_setting = config.CUDA_ENABLE_NRT
+        config.CUDA_ENABLE_NRT = True
+        super().setUp()
+
+    def tearDown(self):
+        config.CUDA_ENABLE_NRT = self.old_nrt_setting
+        super().tearDown()
+
+    def test_list_const_roundtrip(self):
+        lst = [1, 2, 3]
+
+        @cuda.jit
+        def kernel(out):
+            for i in range(len(lst)):
+                out[i] = lst[i]
+
+        out = cuda.to_device(np.zeros(len(lst)))
+
+        kernel[1, 1](out)
+        for g, e in zip(out.copy_to_host(), lst):
+            self.assertEqual(e, g)

numba_cuda/numba/cuda/typing/typeof.py

@@ -193,7 +193,7 @@ def _typeof_list(val, c):
     ty = typeof_impl(val[0], c)
     if ty is None:
         raise ValueError(f"Cannot type list element type {type(val[0])}")
-    return types.List(ty, reflected=True)
+    return ty[::1]
 
 
 @typeof_impl.register(set)