stuartarchibald · Mar 17, 2021 · Mar 16, 2021 · Mar 16, 2021 · Mar 17, 2021 · Mar 17, 2021
diff --git a/docs/source/reference/deprecation.rst b/docs/source/reference/deprecation.rst
@@ -306,3 +306,33 @@ In 0.54.0:
 - ``ptx`` and the inspection methods will always return a dict.
 - Support for indexing into the results of these methods using ``(cc,
   argtypes)`` will be removed.
+
+
+.. _deprecation-strict-strides:
+
+Deprecation of strict strides checking when computing contiguity
+================================================================
+
+The contiguity of device arrays (the ``'C_CONTIGUOUS'`` and ``'F_CONTIGUOUS'``
+elements of the flags of a device array) are computed using relaxed strides
+checking, which matches the default in NumPy since Version 1.12. A config
+variable, :envvar:`NUMBA_NPY_RELAXED_STRIDES_CHECKING`, is provided to force
+computation of these flags using strict strides checking.
+
+This flag is provided to work around any bugs that may be exposed by strict
+strides checking, and will be removed in future.
+
+Schedule
+--------
+
+In 0.54.0:
+
+- Relaxed strides checking will become the default.
+- Strict strides checking will be deprecated.
+
+In 0.55.0:
+
+- Strict strides checking will be removed, if there are no reports of bugs
+  related to relaxed strides checking in 0.54.0 onwards. This plan will be
+  re-examined if bugs related to relaxed strides checking are reported, but may
+  not necessarily change as a result.
diff --git a/docs/source/reference/envvars.rst b/docs/source/reference/envvars.rst
@@ -430,6 +430,21 @@ GPU support
    Setting this variable to 1 also includes the values of arguments to Driver
    API calls in the logs.
 
+.. envvar:: NUMBA_NPY_RELAXED_STRIDES_CHECKING
+
+   By default arrays that inherit from ``numba.misc.dummyarray.Array` (e.g.
+   CUDA and ROCm device arrays) compute their contiguity using relaxed strides
+   checking, which is the default mechanism used by NumPy since version 1.12
+   (see `NPY_RELAXED_STRIDES_CHECKING
+   <https://numpy.org/doc/stable/release/1.8.0-notes.html#npy-relaxed-strides-checking>`_).
+   Setting ``NUMBA_NPY_RELAXED_STRIDES_CHECKING=0`` reverts back to strict
+   strides checking. This option should not normally be needed, but is provided
+   in case it is needed to work around latent bugs related to strict strides
+   checking.
+
+   Strict strides checking is deprecated and may be removed in future. See
+   :ref:`deprecation-strict-strides`.
+
 
 Threading Control
 -----------------

diff --git a/numba/core/config.py b/numba/core/config.py
@@ -337,6 +337,12 @@ def avx_default():
         CUDA_ARRAY_INTERFACE_SYNC = _readenv("NUMBA_CUDA_ARRAY_INTERFACE_SYNC",
                                              int, 1)
 
+        # Compute contiguity of device arrays using the relaxed strides
+        # checking algorithm.
+        NPY_RELAXED_STRIDES_CHECKING = _readenv(
+            "NUMBA_NPY_RELAXED_STRIDES_CHECKING",
+            int, 1)
+
         # HSA Configs
 
         # Disable HSA support

diff --git a/numba/cuda/tests/cudadrv/test_cuda_ndarray.py b/numba/cuda/tests/cudadrv/test_cuda_ndarray.py
@@ -1,8 +1,10 @@
+import itertools
 import numpy as np
 from numba.cuda.cudadrv import devicearray
 from numba import cuda
 from numba.cuda.testing import unittest, CUDATestCase
 from numba.cuda.testing import skip_on_cudasim
+from numba.tests.support import override_config
 
 
 class TestCudaNDArray(CUDATestCase):
@@ -279,33 +281,6 @@ def test_devicearray_contiguous_copy_device(self):
         self.assertTrue(np.all(d.copy_to_host() == a_f))
 
     def test_devicearray_broadcast_host_copy(self):
-        try:
-            broadcast_to = np.broadcast_to
-        except AttributeError:
-            # numpy<1.10 doesn't have broadcast_to. The following implements
-            # a limited broadcast_to that only works along already existing
-            # dimensions of length 1.
-            def broadcast_to(arr, new_shape):
-                new_strides = []
-                for new_length, length, stride in zip(
-                    new_shape, arr.shape, arr.strides
-                ):
-                    if length == 1 and new_length > 1:
-                        new_strides.append(0)
-                    elif new_length == length:
-                        new_strides.append(stride)
-                    else:
-                        raise ValueError(
-                            "cannot broadcast shape {} to shape {}"
-                            .format(arr.shape, new_shape)
-                        )
-                return np.ndarray(
-                    buffer=np.squeeze(arr),
-                    dtype=arr.dtype,
-                    shape=new_shape,
-                    strides=tuple(new_strides),
-                )
-
         broadsize = 4
         coreshape = (2, 3)
         coresize = np.prod(coreshape)
@@ -314,8 +289,8 @@ def broadcast_to(arr, new_shape):
         for dim in range(len(coreshape)):
             newindex = (slice(None),) * dim + (np.newaxis,)
             broadshape = coreshape[:dim] + (broadsize,) + coreshape[dim:]
-            broad_c = broadcast_to(core_c[newindex], broadshape)
-            broad_f = broadcast_to(core_f[newindex], broadshape)
+            broad_c = np.broadcast_to(core_c[newindex], broadshape)
+            broad_f = np.broadcast_to(core_f[newindex], broadshape)
             dbroad_c = cuda.to_device(broad_c)
             dbroad_f = cuda.to_device(broad_f)
             np.testing.assert_array_equal(dbroad_c.copy_to_host(), broad_c)
@@ -343,6 +318,51 @@ def test_devicearray_contiguous_device_strided(self):
             devicearray.errmsg_contiguous_buffer,
             str(e.exception))
 
+    @skip_on_cudasim('DeviceNDArray class not present in simulator')
+    def test_devicearray_relaxed_strides(self):
+        # From the reproducer in Issue #6824.
+
+        with override_config('NPY_RELAXED_STRIDES_CHECKING', 1):
+            # Construct a device array that is contiguous even though
+            # the strides for the first axis (800) are not equal to
+            # the strides * size (10 * 8 = 80) for the previous axis,
+            # because the first axis size is 1.
+            arr = devicearray.DeviceNDArray((1, 10), (800, 8), np.float64)
+
+            # Ensure we still believe the array to be contiguous because
+            # strides checking is relaxed.
+            self.assertTrue(arr.flags['C_CONTIGUOUS'])
+            self.assertTrue(arr.flags['F_CONTIGUOUS'])
+
+    @skip_on_cudasim('DeviceNDArray class not present in simulator')
+    def test_devicearray_strict_strides(self):
+        # From the reproducer in Issue #6824.
+
+        with override_config('NPY_RELAXED_STRIDES_CHECKING', 0):
+            # Construct a device array that is not contiguous because
+            # the strides for the first axis (800) are not equal to
+            # the strides * size (10 * 8 = 80) for the previous axis.
+            arr = devicearray.DeviceNDArray((1, 10), (800, 8), np.float64)
+
+            # Ensure we don't believe the array to be contiguous becase strides
+            # checking is strict.
+            self.assertFalse(arr.flags['C_CONTIGUOUS'])
+            self.assertFalse(arr.flags['F_CONTIGUOUS'])
+
+    def test_c_f_contiguity_matches_numpy(self):
+        # From the reproducer in Issue #4943.
+
+        shapes = ((1, 4), (4, 1))
+        orders = ('C', 'F')
+        with override_config('NPY_RELAXED_STRIDES_CHECKING', 1):
+            for shape, order in itertools.product(shapes, orders):
+                arr = np.ndarray(shape, order=order)
+                d_arr = cuda.to_device(arr)
+                self.assertEqual(arr.flags['C_CONTIGUOUS'],
+                                 d_arr.flags['C_CONTIGUOUS'])
+                self.assertEqual(arr.flags['F_CONTIGUOUS'],
+                                 d_arr.flags['F_CONTIGUOUS'])
+
     @skip_on_cudasim('Typing not done in the simulator')
     def test_devicearray_typing_order_simple_c(self):
         # C-order 1D array

diff --git a/numba/misc/dummyarray.py b/numba/misc/dummyarray.py
@@ -7,6 +7,7 @@
 import numpy as np
 
 from numba import _helperlib
+from numba.core import config
 
 
 Extent = namedtuple("Extent", ["begin", "end"])
@@ -164,7 +165,48 @@ def __init__(self, dims, itemsize):
         self.itemsize = itemsize
         self.size = functools.reduce(operator.mul, self.shape, 1)
         self.extent = self._compute_extent()
-        self.flags = self._compute_layout()
+        if config.NPY_RELAXED_STRIDES_CHECKING:
+            self.flags = self._relaxed_compute_layout()
+        else:
+            self.flags = self._compute_layout()
+
+    def _relaxed_compute_layout(self):
+        # The logic here is based on that in _UpdateContiguousFlags from
+        # numpy/core/src/multiarray/flagsobject.c in NumPy v1.19.1 (commit
+        # 13661ac70).
+        # https://github.com/numpy/numpy/blob/maintenance/1.19.x/numpy/core/src/multiarray/flagsobject.c#L123-L191
+
+        # Records have no dims, and we can treat them as contiguous
+        if not self.dims:
+            return {'C_CONTIGUOUS': True, 'F_CONTIGUOUS': True}
+
+        # If this is a broadcast array then it is not contiguous
+        if any([dim.stride == 0 for dim in self.dims]):
+            return {'C_CONTIGUOUS': False, 'F_CONTIGUOUS': False}
+
+        flags = {'C_CONTIGUOUS': True, 'F_CONTIGUOUS': True}
+
+        # Check C contiguity
+        sd = self.itemsize
+        for dim in reversed(self.dims):
+            if dim.size == 0:
+                # Contiguous by definition
+                return {'C_CONTIGUOUS': True, 'F_CONTIGUOUS': True}
+            if dim.size != 1:
+                if dim.stride != sd:
+                    flags['C_CONTIGUOUS'] = False
+                sd *= dim.size
+
+        # Check F contiguity
+        sd = self.itemsize
+        for dim in self.dims:
+            if dim.size != 1:
+                if dim.stride != sd:
+                    flags['F_CONTIGUOUS'] = False
+                    return flags
+                sd *= dim.size
+
+        return flags
 
     def _compute_layout(self):
         flags = {}

diff --git a/numba/tests/test_dummyarray.py b/numba/tests/test_dummyarray.py
@@ -154,6 +154,13 @@ def test_strided_3d(self):
             self.assertEqual(got.shape, expect.shape)
             self.assertEqual(got.strides, expect.strides)
 
+    def test_issue_2766(self):
+        z = np.empty((1, 2, 3))
+        z = np.transpose(z, axes=(2, 0, 1))
+        arr = Array.from_desc(0, z.shape, z.strides, z.itemsize)
+        self.assertEqual(z.flags['C_CONTIGUOUS'], arr.flags['C_CONTIGUOUS'])
+        self.assertEqual(z.flags['F_CONTIGUOUS'], arr.flags['F_CONTIGUOUS'])
+
 
 class TestReshape(unittest.TestCase):
     def test_reshape_2d2d(self):