test: fix broken simulator interface

Author: cpcloud

numba_cuda/numba/cuda/simulator/cudadrv/devicearray.py

@@ -8,6 +8,8 @@
 
 from contextlib import contextmanager
 from numba.cuda.np.numpy_support import numpy_version
+from numba.cuda.np import numpy_support
+from numba.cuda import types
 
 import numpy as np
 
@@ -57,8 +59,8 @@ def __wrap_if_fake(self, item):
             return item
 
     def __getattr__(self, attrname):
-        if attrname in dir(self._item._ary):  # For e.g. array size.
-            return self.__wrap_if_fake(getattr(self._item._ary, attrname))
+        if (value := getattr(self._item, attrname, None)) is not None:
+            return self.__wrap_if_fake(value)
         else:
             return self.__wrap_if_fake(self._item.__getitem__(attrname))
 
@@ -109,6 +111,23 @@ def __init__(self, ary, stream=0):
         self._ary = ary
         self.stream = stream
 
+    @property
+    def _numba_type_(self):
+        """
+        Magic attribute expected by Numba to get the numba type that
+        represents this object.
+        """
+        broadcast = 0 in self.strides
+        if self.is_c_contiguous() and not broadcast:
+            layout = "C"
+        elif self.is_f_contiguous() and not broadcast:
+            layout = "F"
+        else:
+            layout = "A"
+
+        dtype = numpy_support.from_dtype(self._ary.dtype)
+        return types.Array(dtype, self._ary.ndim, layout)
+
     @property
     def alloc_size(self):
         return self._ary.nbytes

numba_cuda/numba/cuda/testing.py

@@ -191,9 +191,9 @@ def skip_on_cudasim(reason):
     return unittest.skipIf(config.ENABLE_CUDASIM, reason)
 
 
-def skip_on_standalone_numba_cuda(reason):
-    """Skip this test if running on standalone numba_cuda"""
-    return unittest.skipIf(not HAS_NUMBA, reason)
+skip_on_standalone_numba_cuda = unittest.skipUnless(
+    HAS_NUMBA, "requires base numba install"
+)
 
 
 def skip_unless_cudasim(reason):

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

@@ -5,7 +5,7 @@
 from textwrap import dedent
 
 from numba import cuda
-from numba.cuda import uint32, uint64, float32, float64
+from numba.cuda import uint32, uint64, float32, float64, int32
 from numba.cuda.testing import unittest, CUDATestCase, cc_X_or_above
 from numba.cuda.core import config
 
@@ -239,19 +239,19 @@ def atomic_add_double_3(ary):
 
 def atomic_sub(ary):
     atomic_binary_1dim_shared(
-        ary, ary, 1, uint32, 32, cuda.atomic.sub, atomic_cast_none, 0, False
+        ary, ary, 1, int32, 32, cuda.atomic.sub, atomic_cast_none, 0, False
     )
 
 
 def atomic_sub2(ary):
     atomic_binary_2dim_shared(
-        ary, 1, uint32, (4, 8), cuda.atomic.sub, atomic_cast_none, False
+        ary, 1, int32, (4, 8), cuda.atomic.sub, atomic_cast_none, False
     )
 
 
 def atomic_sub3(ary):
     atomic_binary_2dim_shared(
-        ary, 1, uint32, (4, 8), cuda.atomic.sub, atomic_cast_to_uint64, False
+        ary, 1, int32, (4, 8), cuda.atomic.sub, atomic_cast_to_uint64, False
     )
 
 
@@ -789,7 +789,7 @@ def test_atomic_add_double_global_3(self):
         self.assertCorrectFloat64Atomics(cuda_func, shared=False)
 
     def test_atomic_sub(self):
-        ary = np.random.randint(0, 32, size=32).astype(np.int32)
+        ary = np.random.randint(0, 32, size=32, dtype=np.int32)
         orig = ary.copy()
         cuda_atomic_sub = cuda.jit("void(int32[:])")(atomic_sub)
         cuda_atomic_sub[1, 32](ary)
@@ -801,16 +801,16 @@ def test_atomic_sub(self):
         self.assertTrue(np.all(ary == gold))
 
     def test_atomic_sub2(self):
-        ary = np.random.randint(0, 32, size=32).astype(np.uint32).reshape(4, 8)
+        ary = np.random.randint(0, 32, size=(4, 8), dtype=np.int32)
         orig = ary.copy()
-        cuda_atomic_sub2 = cuda.jit("void(uint32[:,:])")(atomic_sub2)
+        cuda_atomic_sub2 = cuda.jit("void(int32[:,:])")(atomic_sub2)
         cuda_atomic_sub2[1, (4, 8)](ary)
         self.assertTrue(np.all(ary == orig - 1))
 
     def test_atomic_sub3(self):
-        ary = np.random.randint(0, 32, size=32).astype(np.uint32).reshape(4, 8)
+        ary = np.random.randint(0, 32, size=(4, 8), dtype=np.uint32)
         orig = ary.copy()
-        cuda_atomic_sub3 = cuda.jit("void(uint32[:,:])")(atomic_sub3)
+        cuda_atomic_sub3 = cuda.jit("void(int32[:,:])")(atomic_sub3)
         cuda_atomic_sub3[1, (4, 8)](ary)
         self.assertTrue(np.all(ary == orig - 1))
 

testing/pytest.ini

@@ -23,3 +23,4 @@ filterwarnings =
     ignore:.*Compilation is falling back to object mode WITHOUT looplifting enabled.*:numba.cuda.core.errors.NumbaWarning
     ignore:\nCompilation is falling back to object mode WITHOUT looplifting enabled.*:numba.core.errors.NumbaWarning
     ignore:overflow encountered in scalar .+:RuntimeWarning
+    ignore:.*Host array used in CUDA kernel will incur copy overhead.*:numba.cuda.core.errors.NumbaPerformanceWarning