Vendor numpy median, percentile, and quantile reduction tests

numba_cuda/numba/cuda/np/arrayobj.py

@@ -5504,37 +5504,31 @@ def _array_copy(context, builder, sig, args):
     dest_data = ret.data
 
     assert rettype.layout in "CF"
-    if arytype.layout == rettype.layout:
-        # Fast path: memcpy
-        cgutils.raw_memcpy(
-            builder, dest_data, src_data, ary.nitems, ary.itemsize, align=1
-        )
 
-    else:
-        src_strides = cgutils.unpack_tuple(builder, ary.strides)
-        dest_strides = cgutils.unpack_tuple(builder, ret.strides)
-        intp_t = context.get_value_type(types.intp)
+    src_strides = cgutils.unpack_tuple(builder, ary.strides)
+    dest_strides = cgutils.unpack_tuple(builder, ret.strides)
+    intp_t = context.get_value_type(types.intp)
 
-        with cgutils.loop_nest(builder, shapes, intp_t) as indices:
-            src_ptr = cgutils.get_item_pointer2(
-                context,
-                builder,
-                src_data,
-                shapes,
-                src_strides,
-                arytype.layout,
-                indices,
-            )
-            dest_ptr = cgutils.get_item_pointer2(
-                context,
-                builder,
-                dest_data,
-                shapes,
-                dest_strides,
-                rettype.layout,
-                indices,
-            )
-            builder.store(builder.load(src_ptr), dest_ptr)
+    with cgutils.loop_nest(builder, shapes, intp_t) as indices:
+        src_ptr = cgutils.get_item_pointer2(
+            context,
+            builder,
+            src_data,
+            shapes,
+            src_strides,
+            arytype.layout,
+            indices,
+        )
+        dest_ptr = cgutils.get_item_pointer2(
+            context,
+            builder,
+            dest_data,
+            shapes,
+            dest_strides,
+            rettype.layout,
+            indices,
+        )
+        builder.store(builder.load(src_ptr), dest_ptr)
 
     return impl_ret_new_ref(context, builder, sig.return_type, ret._getvalue())
 

numba_cuda/numba/cuda/target.py

@@ -275,6 +275,34 @@ 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):
+        import numpy as np
+
+        constvals = [
+            self.get_constant(listty.dtype, i) for i in iter(np.array(lst))
+        ]
+        instance = self.build_list(builder, listty, constvals)
+        # create constant address space version of the list
+        lmod = builder.module
+
+        constlistty = instance.type
+        constlist = ir.Constant(constlistty, instance)
+        addrspace = nvvm.ADDRSPACE_CONSTANT
+        gv = cgutils.add_global_variable(
+            lmod, constlist.type, "_cudapy_clist", addrspace=addrspace
+        )
+        gv.linkage = "internal"
+        gv.global_constant = True
+        gv.initializer = constlist
+
+        # Convert to generic address-space
+        ptrty = ir.PointerType(constlistty)
+        genptr = builder.addrspacecast(gv, ptrty, "generic")
+        lst = cgutils.create_struct_proxy(listty)(
+            self, builder, value=builder.load(genptr)
+        )
+        return lst._getvalue()
+
     def make_constant_array(self, builder, aryty, arr):
         """
         Unlike the parent version.  This returns a a pointer in the constant

numba_cuda/numba/cuda/testing.py

@@ -185,6 +185,17 @@ def assertFileCheckMatches(
             )
 
 
+class NRTEnablingCUDATestCase(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 skip_on_cudasim(reason):
     """Skip this test if running on the CUDA simulator"""
     assert isinstance(reason, str)

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

@@ -3,9 +3,8 @@
 
 import numpy as np
 from numba import cuda
-from numba.cuda.testing import CUDATestCase
+from numba.cuda.testing import NRTEnablingCUDATestCase
 import unittest
-from numba.cuda import config
 
 
 def reinterpret_array_type(byte_arr, start, stop, output):
@@ -14,16 +13,7 @@ def reinterpret_array_type(byte_arr, start, stop, output):
     output[0] = val
 
 
-class TestCudaArrayMethods(CUDATestCase):
-    def setUp(self):
-        self.old_nrt_setting = config.CUDA_ENABLE_NRT
-        config.CUDA_ENABLE_NRT = True
-        super(TestCudaArrayMethods, self).setUp()
-
-    def tearDown(self):
-        config.CUDA_ENABLE_NRT = self.old_nrt_setting
-        super(TestCudaArrayMethods, self).tearDown()
-
+class TestCudaArrayMethods(NRTEnablingCUDATestCase):
     def test_reinterpret_array_type(self):
         """
         Reinterpret byte array as int32 in the GPU.