diff --git a/numba_cuda/numba/cuda/np/arrayobj.py b/numba_cuda/numba/cuda/np/arrayobj.py index de8834e91..52184a0a2 100644 --- a/numba_cuda/numba/cuda/np/arrayobj.py +++ b/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) diff --git a/numba_cuda/numba/cuda/target.py b/numba_cuda/numba/cuda/target.py index dde6ab760..0f6efd64a 100644 --- a/numba_cuda/numba/cuda/target.py +++ b/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) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_lists.py b/numba_cuda/numba/cuda/tests/cudapy/test_lists.py new file mode 100644 index 000000000..71a5de4a8 --- /dev/null +++ b/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) diff --git a/numba_cuda/numba/cuda/typing/typeof.py b/numba_cuda/numba/cuda/typing/typeof.py index a3091d282..a5c4b128d 100644 --- a/numba_cuda/numba/cuda/typing/typeof.py +++ b/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)