From ef8dfae81793a7fd385b329a2604cacd9e0448a5 Mon Sep 17 00:00:00 2001 From: brandon-b-miller Date: Wed, 5 Nov 2025 16:00:25 -0800 Subject: [PATCH 1/4] initial --- numba_cuda/numba/cuda/np/arrayobj.py | 8 +++++ numba_cuda/numba/cuda/target.py | 51 ++++++++++++++++++++++++++-- 2 files changed, 57 insertions(+), 2 deletions(-) diff --git a/numba_cuda/numba/cuda/np/arrayobj.py b/numba_cuda/numba/cuda/np/arrayobj.py index de8834e91..2176fe740 100644 --- a/numba_cuda/numba/cuda/np/arrayobj.py +++ b/numba_cuda/numba/cuda/np/arrayobj.py @@ -3656,6 +3656,14 @@ def constant_record(context, builder, ty, pyval): return cgutils.alloca_once_value(builder, val) +@lower_constant(types.List) +def constant_list(context, builder, ty, pyval): + """ + Create a constant list (mechanism is target-dependent). + """ + return context.make_constant_list(builder, ty, pyval) + + @lower_constant(types.Bytes) def constant_bytes(context, builder, ty, pyval): """ diff --git a/numba_cuda/numba/cuda/target.py b/numba_cuda/numba/cuda/target.py index dde6ab760..17f3efb49 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) @@ -274,6 +273,54 @@ def mangler(self, name, argtypes, *, abi_tags=(), uid=None): name, argtypes, abi_tags=abi_tags, uid=uid ) + def make_constant_list(self, builder, listty, lst): + """ + Create a list structure with data in constant address space. + The list metadata is in normal memory but the data array is in constant memory. + """ + import numpy as np + from numba.cuda.cpython.listobj import ListInstance + + lmod = builder.module + + # Convert to array and serialize + arr = np.array(lst) + nitems = len(arr) + + constvals = [ + self.get_constant(types.byte, i) + for i in iter(arr.tobytes(order="C")) + ] + constaryty = ir.ArrayType(ir.IntType(8), len(constvals)) + constary = ir.Constant(constaryty, constvals) + + # Create constant memory data + addrspace = nvvm.ADDRSPACE_CONSTANT + gv = cgutils.add_global_variable( + lmod, constary.type, "_cudapy_clist", addrspace=addrspace + ) + gv.linkage = "internal" + gv.global_constant = True + gv.initializer = constary + + lldtype = self.get_data_type(listty.dtype) + align = self.get_abi_sizeof(lldtype) + gv.align = 2 ** (align - 1).bit_length() + + ptrty = ir.PointerType(ir.IntType(8)) + genptr = builder.addrspacecast(gv, ptrty, "generic") + + list_inst = ListInstance.allocate(self, builder, listty, nitems) + data_ptrty = self.get_data_type(listty.dtype).as_pointer() + const_data_ptr = builder.bitcast(genptr, data_ptrty) + payload = list_inst._payload + + # replace with constant mem + payload.data = builder.ptrtoint(const_data_ptr, payload.data.type) + list_inst.size = self.get_constant(types.intp, nitems) + + return list_inst.value + def make_constant_array(self, builder, aryty, arr): """ Unlike the parent version. This returns a a pointer in the constant From 956c3ca14d15d7ce0ef94353de9154f88dea8871 Mon Sep 17 00:00:00 2001 From: brandon-b-miller Date: Thu, 6 Nov 2025 06:45:47 -0800 Subject: [PATCH 2/4] tests --- .../numba/cuda/tests/cudapy/test_lists.py | 38 +++++++++++++++++++ 1 file changed, 38 insertions(+) create mode 100644 numba_cuda/numba/cuda/tests/cudapy/test_lists.py 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..bcb68e0af --- /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_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) From 6775a692a92155d0b29beff57ed76eee60cc288f Mon Sep 17 00:00:00 2001 From: brandon-b-miller Date: Thu, 6 Nov 2025 06:58:00 -0800 Subject: [PATCH 3/4] small changes --- numba_cuda/numba/cuda/target.py | 12 ++++++++---- numba_cuda/numba/cuda/tests/cudapy/test_lists.py | 2 +- 2 files changed, 9 insertions(+), 5 deletions(-) diff --git a/numba_cuda/numba/cuda/target.py b/numba_cuda/numba/cuda/target.py index 17f3efb49..91ed9fb36 100644 --- a/numba_cuda/numba/cuda/target.py +++ b/numba_cuda/numba/cuda/target.py @@ -292,22 +292,26 @@ def make_constant_list(self, builder, listty, lst): for i in iter(arr.tobytes(order="C")) ] constaryty = ir.ArrayType(ir.IntType(8), len(constvals)) - constary = ir.Constant(constaryty, constvals) + constary = ir.Constant( + constaryty, constvals + ) # constant llvm value with the values - # Create constant memory data + # create a global variable addrspace = nvvm.ADDRSPACE_CONSTANT gv = cgutils.add_global_variable( lmod, constary.type, "_cudapy_clist", addrspace=addrspace ) gv.linkage = "internal" gv.global_constant = True - gv.initializer = constary + gv.initializer = ( + constary # put the gobal variable initializer to the constant array + ) lldtype = self.get_data_type(listty.dtype) align = self.get_abi_sizeof(lldtype) gv.align = 2 ** (align - 1).bit_length() - ptrty = ir.PointerType(ir.IntType(8)) + ptrty = ir.PointerType(ir.IntType(8)) # an int8 pointer genptr = builder.addrspacecast(gv, ptrty, "generic") list_inst = ListInstance.allocate(self, builder, listty, nitems) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_lists.py b/numba_cuda/numba/cuda/tests/cudapy/test_lists.py index bcb68e0af..71a5de4a8 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_lists.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_lists.py @@ -23,7 +23,7 @@ def tearDown(self): config.CUDA_ENABLE_NRT = self.old_nrt_setting super().tearDown() - def test_list_roundtrip(self): + def test_list_const_roundtrip(self): lst = [1, 2, 3] @cuda.jit From d843267590774b51a052f96c621991662a1afc04 Mon Sep 17 00:00:00 2001 From: brandon-b-miller Date: Thu, 6 Nov 2025 07:22:28 -0800 Subject: [PATCH 4/4] suggestion from graham Co-authored-by: Graham Markall --- numba_cuda/numba/cuda/np/arrayobj.py | 10 +---- numba_cuda/numba/cuda/target.py | 52 -------------------------- numba_cuda/numba/cuda/typing/typeof.py | 2 +- 3 files changed, 3 insertions(+), 61 deletions(-) diff --git a/numba_cuda/numba/cuda/np/arrayobj.py b/numba_cuda/numba/cuda/np/arrayobj.py index 2176fe740..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) @@ -3656,14 +3658,6 @@ def constant_record(context, builder, ty, pyval): return cgutils.alloca_once_value(builder, val) -@lower_constant(types.List) -def constant_list(context, builder, ty, pyval): - """ - Create a constant list (mechanism is target-dependent). - """ - return context.make_constant_list(builder, ty, pyval) - - @lower_constant(types.Bytes) def constant_bytes(context, builder, ty, pyval): """ diff --git a/numba_cuda/numba/cuda/target.py b/numba_cuda/numba/cuda/target.py index 91ed9fb36..0f6efd64a 100644 --- a/numba_cuda/numba/cuda/target.py +++ b/numba_cuda/numba/cuda/target.py @@ -273,58 +273,6 @@ def mangler(self, name, argtypes, *, abi_tags=(), uid=None): name, argtypes, abi_tags=abi_tags, uid=uid ) - def make_constant_list(self, builder, listty, lst): - """ - Create a list structure with data in constant address space. - The list metadata is in normal memory but the data array is in constant memory. - """ - import numpy as np - from numba.cuda.cpython.listobj import ListInstance - - lmod = builder.module - - # Convert to array and serialize - arr = np.array(lst) - nitems = len(arr) - - constvals = [ - self.get_constant(types.byte, i) - for i in iter(arr.tobytes(order="C")) - ] - constaryty = ir.ArrayType(ir.IntType(8), len(constvals)) - constary = ir.Constant( - constaryty, constvals - ) # constant llvm value with the values - - # create a global variable - addrspace = nvvm.ADDRSPACE_CONSTANT - gv = cgutils.add_global_variable( - lmod, constary.type, "_cudapy_clist", addrspace=addrspace - ) - gv.linkage = "internal" - gv.global_constant = True - gv.initializer = ( - constary # put the gobal variable initializer to the constant array - ) - - lldtype = self.get_data_type(listty.dtype) - align = self.get_abi_sizeof(lldtype) - gv.align = 2 ** (align - 1).bit_length() - - ptrty = ir.PointerType(ir.IntType(8)) # an int8 pointer - genptr = builder.addrspacecast(gv, ptrty, "generic") - - list_inst = ListInstance.allocate(self, builder, listty, nitems) - data_ptrty = self.get_data_type(listty.dtype).as_pointer() - const_data_ptr = builder.bitcast(genptr, data_ptrty) - payload = list_inst._payload - - # replace with constant mem - payload.data = builder.ptrtoint(const_data_ptr, payload.data.type) - list_inst.size = self.get_constant(types.intp, nitems) - - return list_inst.value - def make_constant_array(self, builder, aryty, arr): """ Unlike the parent version. This returns a a pointer in the constant 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)