suggestion from graham

Co-authored-by: Graham Markall <[email protected]>

Author: brandon-b-miller

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):
     """

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

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)