Repository: numba/numba

Issue:
Slicing device array with negative step produces wrong result
## Reporting a bug

- [X] I am using the latest released version of Numba (most recent is visible in
 the change log (https://github.com/numba/numba/blob/master/CHANGE_LOG).
- [X] I have included below a minimal working reproducer (if you are unsure how
 to write one see http://matthewrocklin.com/blog/work/2018/02/28/minimal-bug-reports).

Consider
```python
import numpy as np
import numba
import numba.cuda as cuda
arr = np.arange(4)
d_arr = cuda.to_device(arr)

print('numba version', numba.__version__)
print('numpy array:',arr[::-1].__array_interface__)
print('numba.cuda device array', d_arr[::-1].__cuda_array_interface__)
```
that outputs
```
numba version 0.42.0
numpy array: {'data': (94375105644824, False), 'strides': (-8,), 'descr': [('', '<i8')], 'typestr': '<i8', 'shape': (4,), 'version': 3}
numba.cuda device array {'shape': (-3,), 'strides': (-8,), 'data': (None, False), 'typestr': '<i8', 'version': 0}
```
Clearly, the shape `(-3,)` of `d_arr[::-1]` is wrong, it should be `(4,)`.

As a result, `d_arr[::-1].copy_to_host()` fails:
```python-traceback
Traceback (most recent call last):
  File "pyarrow/tests/numba_bug.py", line 9, in <module>
    d_arr[::-1].copy_to_host()
  File "/home/pearu/miniconda3/envs/pyarrow-dev/lib/python3.7/site-packages/numba/cuda/cudadrv/devices.py", line 212, in _require_cuda_context
    return fn(*args, **kws)
  File "/home/pearu/miniconda3/envs/pyarrow-dev/lib/python3.7/site-packages/numba/cuda/cudadrv/devicearray.py", line 240, in copy_to_host
    strides=self.strides, buffer=hostary)
ValueError: strides is incompatible with shape of requested array and size of buffer
```

Discussion:

Patch:
diff --git a/numba/cuda/cudadrv/devicearray.py b/numba/cuda/cudadrv/devicearray.py
index c5503b395..1f30862f3 100644
--- a/numba/cuda/cudadrv/devicearray.py
+++ b/numba/cuda/cudadrv/devicearray.py
@@ -220,6 +220,10 @@ class DeviceNDArrayBase(object):
 
             result_array = d_arr.copy_to_host()
         """
+        if any(s < 0 for s in self.strides):
+            msg = 'D->H copy not implemented for negative strides: {}'
+            raise NotImplementedError(msg.format(self.strides))
+        assert self.alloc_size >= 0, "Negative memory size"
         stream = self._default_stream(stream)
         if ary is None:
             hostary = np.empty(shape=self.alloc_size, dtype=np.byte)
@@ -227,7 +231,6 @@ class DeviceNDArrayBase(object):
             check_array_compatibility(self, ary)
             hostary = ary
 
-        assert self.alloc_size >= 0, "Negative memory size"
         if self.alloc_size != 0:
             _driver.device_to_host(hostary, self, self.alloc_size, stream=stream)
 
diff --git a/numba/cuda/cudadrv/driver.py b/numba/cuda/cudadrv/driver.py
index 753420c60..0a07dd4be 100644
--- a/numba/cuda/cudadrv/driver.py
+++ b/numba/cuda/cudadrv/driver.py
@@ -1715,7 +1715,7 @@ def device_memory_size(devmem):
         s, e = device_extents(devmem)
         sz = e - s
         devmem._cuda_memsize_ = sz
-    assert sz > 0, "zero length array"
+    assert sz >= 0, "{} length array".format(sz)
     return sz
 
 
diff --git a/numba/dummyarray.py b/numba/dummyarray.py
index f04b9d490..3a69c1bd9 100644
--- a/numba/dummyarray.py
+++ b/numba/dummyarray.py
@@ -43,8 +43,6 @@ class Dim(object):
     __slots__ = 'start', 'stop', 'size', 'stride', 'single'
 
     def __init__(self, start, stop, size, stride, single):
-        if stop < start:
-            raise ValueError("end offset is before start offset")
         self.start = start
         self.stop = stop
         self.size = size
@@ -54,54 +52,31 @@ class Dim(object):
 
     def __getitem__(self, item):
         if isinstance(item, slice):
-            start, stop, step = item.start, item.stop, item.step
-            single = False
-        else:
-            single = True
-            start = item
-            stop = start + 1
-            step = None
-
-        # Default values
-        #   Start value is default to zero
-        if start is None:
-            start = 0
-        #   Stop value is default to self.size
-        if stop is None:
-            stop = self.size
-        #   Step is default to 1
-        if step is None:
-            step = 1
-
-        stride = step * self.stride
-
-        # Compute start in bytes
-        if start >= 0:
-            start = self.start + start * self.stride
-        else:
-            start = self.stop + start * self.stride
-        start = max(start, self.start)
-
-        # Compute stop in bytes
-        if stop >= 0:
-            stop = self.start + stop * self.stride
-        else:
-            stop = self.stop + stop * self.stride
-        stop = min(stop, self.stop)
-
-        # Clip stop
-        if (stop - start) > self.size * self.stride:
-            stop = start + self.size * stride
-
-        if stop < start:
-            start = stop
-            size = 0
-        elif stride == 0:
-            size = 1 if single else ((stop - start) // step)
+            start, stop, step = item.indices(self.size)
+            stride = step * self.stride
+            start = self.start + start * abs(self.stride)
+            stop = self.start + stop * abs(self.stride)
+            if stride == 0:
+                size = 1
+            else:
+                size = _compute_size(start, stop, stride)
+            ret = Dim(
+                start=start,
+                stop=stop,
+                size=size,
+                stride=stride,
+                single=False
+            )
+            return ret
         else:
-            size = (stop - start + (stride - 1)) // stride
-
-        return Dim(start, stop, size, stride, single)
+            sliced = self[item:item + 1]
+            return Dim(
+                start=sliced.start,
+                stop=sliced.stop,
+                size=sliced.size,
+                stride=sliced.stride,
+                single=True,
+            )
 
     def get_offset(self, idx):
         return self.start + idx * self.stride
@@ -428,3 +403,17 @@ def is_element_indexing(item, ndim):
 
     return False
 
+
+def _compute_size(start, stop, step):
+    """Algorithm adapted from cpython rangeobject.c
+    """
+    if step > 0:
+        lo = start
+        hi = stop
+    else:
+        lo = stop
+        hi = start
+        step = -step
+    if lo >= hi:
+        return 0
+    return (hi - lo - 1) // step + 1
diff --git a/numba/ir_utils.py b/numba/ir_utils.py
index 33de75d91..51c6886f8 100644
--- a/numba/ir_utils.py
+++ b/numba/ir_utils.py
@@ -256,6 +256,25 @@ def mk_loop_header(typemap, phi_var, calltypes, scope, loc):
     return header_block
 
 
+def find_op_typ(op, arg_typs):
+    for ft in typing.templates.builtin_registry.functions:
+        if ft.key == op:
+            try:
+                func_typ = types.Function(ft).get_call_type(typing.Context(),
+                                                            arg_typs, {})
+            except TypingError:
+                func_typ = None
+            if func_typ is not None:
+                return func_typ
+    else:
+        for f, ft in typing.templates.builtin_registry.globals:
+            if f == op:
+                return ft.get_call_type(typing.Context(),
+                                                    arg_typs, {})
+
+    raise RuntimeError("unknown array operation")
+
+
 def legalize_names(varnames):
     """returns a dictionary for conversion of variable names to legal
     parameter names.
diff --git a/numba/parfor.py b/numba/parfor.py
index 29251ca39..325d559c5 100644
--- a/numba/parfor.py
+++ b/numba/parfor.py
@@ -41,6 +41,7 @@ from numba.ir_utils import (
     get_np_ufunc_typ,
     mk_range_block,
     mk_loop_header,
+    find_op_typ,
     get_name_var_table,
     replace_vars,
     replace_vars_inner,
@@ -2536,13 +2537,13 @@ def _arrayexpr_tree_to_ir(
             el_typ1 = typemap[arg_vars[0].name]
             if len(arg_vars) == 2:
                 el_typ2 = typemap[arg_vars[1].name]
-                func_typ = typingctx.resolve_function_type(op, (el_typ1, el_typ), {})
+                func_typ = find_op_typ(op, [el_typ1, el_typ2])
                 ir_expr = ir.Expr.binop(op, arg_vars[0], arg_vars[1], loc)
                 if op == operator.truediv:
                     func_typ, ir_expr = _gen_np_divide(
                         arg_vars[0], arg_vars[1], out_ir, typemap)
             else:
-                func_typ = typingctx.resolve_function_type(op, (el_typ1,), {})
+                func_typ = find_op_typ(op, [el_typ1])
                 ir_expr = ir.Expr.unary(op, arg_vars[0], loc)
             calltypes[ir_expr] = func_typ
             el_typ = func_typ.return_type
@@ -3728,7 +3729,7 @@ def apply_copies_parfor(parfor, var_dict, name_var_table,
 ir_utils.apply_copy_propagate_extensions[Parfor] = apply_copies_parfor
 
 
-def push_call_vars(blocks, saved_globals, saved_getattrs, nested=False):
+def push_call_vars(blocks, saved_globals, saved_getattrs):
     """push call variables to right before their call site.
     assuming one global/getattr is created for each call site and control flow
     doesn't change it.
@@ -3752,11 +3753,11 @@ def push_call_vars(blocks, saved_globals, saved_getattrs, nested=False):
                             saved_getattrs[lhs.name] = stmt
                             block_defs.add(lhs.name)
 
-            if not nested and isinstance(stmt, Parfor):
+            if isinstance(stmt, Parfor):
                 for s in stmt.init_block.body:
                     process_assign(s)
                 pblocks = stmt.loop_body.copy()
-                push_call_vars(pblocks, saved_globals, saved_getattrs, nested=True)
+                push_call_vars(pblocks, saved_globals, saved_getattrs)
                 new_body.append(stmt)
                 continue
             else:
diff --git a/numba/stencilparfor.py b/numba/stencilparfor.py
index 4498162e3..f5368aaac 100644
--- a/numba/stencilparfor.py
+++ b/numba/stencilparfor.py
@@ -602,8 +602,8 @@ class StencilPass(object):
             self.typemap[index_var.name] = types.intp
             index_call = ir.Expr.binop(operator.add, old_index_var,
                                                 offset_var, loc)
-            self.calltypes[index_call] = self.typingctx.resolve_function_type(
-                                         operator.add, (types.intp, types.intp), {})
+            self.calltypes[index_call] = ir_utils.find_op_typ(operator.add,
+                                        [types.intp, types.intp])
             index_assign = ir.Assign(index_call, index_var, loc)
             out_nodes.append(index_assign)
             index_vars.append(index_var)
diff --git a/numba/targets/arraymath.py b/numba/targets/arraymath.py
index 87fdc4339..26e72cb73 100644
--- a/numba/targets/arraymath.py
+++ b/numba/targets/arraymath.py
@@ -7,7 +7,6 @@ from __future__ import print_function, absolute_import, division
 import math
 from collections import namedtuple
 from enum import IntEnum
-from functools import partial
 
 import numpy as np
 
@@ -352,18 +351,10 @@ def array_std(context, builder, sig, args):
     return impl_ret_untracked(context, builder, sig.return_type, res)
 
 
-def zero_dim_msg(fn_name):
-    msg = ("zero-size array to reduction operation "
-           "{0} which has no identity".format(fn_name))
-    return msg
-
-
 @lower_builtin(np.min, types.Array)
 @lower_builtin("array.min", types.Array)
 def array_min(context, builder, sig, args):
     ty = sig.args[0].dtype
-    MSG = zero_dim_msg('minimum')
-
     if isinstance(ty, (types.NPDatetime, types.NPTimedelta)):
         # NaT is smaller than every other value, but it is
         # ignored as far as min() is concerned.
@@ -371,8 +362,8 @@ def array_min(context, builder, sig, args):
 
         def array_min_impl(arry):
             if arry.size == 0:
-                raise ValueError(MSG)
-
+                raise ValueError(("zero-size array to reduction operation "
+                                  "minimum which has no identity"))
             min_value = nat
             it = np.nditer(arry)
             for view in it:
@@ -387,37 +378,21 @@ def array_min(context, builder, sig, args):
                     min_value = v
             return min_value
 
-    elif isinstance(ty, types.Complex):
-        def array_min_impl(arry):
-            if arry.size == 0:
-                raise ValueError(MSG)
-
-            it = np.nditer(arry)
-            min_value = next(it).take(0)
-
-            for view in it:
-                v = view.item()
-                if v.real < min_value.real:
-                    min_value = v
-                elif v.real == min_value.real:
-                    if v.imag < min_value.imag:
-                        min_value = v
-            return min_value
-
     else:
         def array_min_impl(arry):
             if arry.size == 0:
-                raise ValueError(MSG)
-
+                raise ValueError(("zero-size array to reduction operation "
+                                  "minimum which has no identity"))
             it = np.nditer(arry)
-            min_value = next(it).take(0)
+            for view in it:
+                min_value = view.item()
+                break
 
             for view in it:
                 v = view.item()
                 if v < min_value:
                     min_value = v
             return min_value
-
     res = context.compile_internal(builder, array_min_impl, sig, args)
     return impl_ret_borrowed(context, builder, sig.return_type, res)
 
@@ -425,39 +400,20 @@ def array_min(context, builder, sig, args):
 @lower_builtin(np.max, types.Array)
 @lower_builtin("array.max", types.Array)
 def array_max(context, builder, sig, args):
-    ty = sig.args[0].dtype
-    MSG = zero_dim_msg('maximum')
-
-    if isinstance(ty, types.Complex):
-        def array_max_impl(arry):
-            if arry.size == 0:
-                raise ValueError(MSG)
-
-            it = np.nditer(arry)
-            max_value = next(it).take(0)
-
-            for view in it:
-                v = view.item()
-                if v.real > max_value.real:
-                    max_value = v
-                elif v.real == max_value.real:
-                    if v.imag > max_value.imag:
-                        max_value = v
-            return max_value
-    else:
-        def array_max_impl(arry):
-            if arry.size == 0:
-                raise ValueError(MSG)
-
-            it = np.nditer(arry)
-            max_value = next(it).take(0)
-
-            for view in it:
-                v = view.item()
-                if v > max_value:
-                    max_value = v
-            return max_value
+    def array_max_impl(arry):
+        if arry.size == 0:
+            raise ValueError(("zero-size array to reduction operation "
+                                "maximum which has no identity"))
+        it = np.nditer(arry)
+        for view in it:
+            max_value = view.item()
+            break
 
+        for view in it:
+            v = view.item()
+            if v > max_value:
+                max_value = v
+        return max_value
     res = context.compile_internal(builder, array_max_impl, sig, args)
     return impl_ret_borrowed(context, builder, sig.return_type, res)
 
@@ -574,81 +530,46 @@ def get_isnan(dtype):
             return False
         return _trivial_isnan
 
-@register_jitable
-def less_than(a, b):
-    return a < b
-
-@register_jitable
-def greater_than(a, b):
-    return a > b
-
-@register_jitable
-def check_array(a):
-    if a.size == 0:
-        raise ValueError('zero-size array to reduction operation not possible')
-
-def nan_min_max_factory(comparison_op, is_complex_dtype):
-
-    if is_complex_dtype:
-        def impl(a):
-            arr = np.asarray(a)
-            check_array(arr)
-            it = np.nditer(arr)
-            return_val = next(it).take(0)
-            for view in it:
-                v = view.item()
-                if np.isnan(return_val.real) and not np.isnan(v.real):
-                    return_val = v
-                else:
-                    if comparison_op(v.real, return_val.real):
-                        return_val = v
-                    elif v.real == return_val.real:
-                        if comparison_op(v.imag, return_val.imag):
-                            return_val = v
-            return return_val
-    else:
-        def impl(a):
-            arr = np.asarray(a)
-            check_array(arr)
-            it = np.nditer(arr)
-            return_val = next(it).take(0)
-            for view in it:
-                v = view.item()
-                if not np.isnan(v):
-                    if not comparison_op(return_val, v):
-                        return_val = v
-            return return_val
-
-    return impl
-
-real_nanmin = register_jitable(
-    nan_min_max_factory(less_than, is_complex_dtype=False)
-)
-real_nanmax = register_jitable(
-    nan_min_max_factory(greater_than, is_complex_dtype=False)
-)
-complex_nanmin = register_jitable(
-    nan_min_max_factory(less_than, is_complex_dtype=True)
-)
-complex_nanmax = register_jitable(
-    nan_min_max_factory(greater_than, is_complex_dtype=True)
-)
 
 @overload(np.nanmin)
 def np_nanmin(a):
-    dt = determine_dtype(a)
-    if np.issubdtype(dt, np.complexfloating):
-        return complex_nanmin
-    else:
-        return real_nanmin
+    if not isinstance(a, types.Array):
+        return
+    isnan = get_isnan(a.dtype)
+
+    def nanmin_impl(a):
+        if a.size == 0:
+            raise ValueError("nanmin(): empty array")
+        for view in np.nditer(a):
+            minval = view.item()
+            break
+        for view in np.nditer(a):
+            v = view.item()
+            if not minval < v and not isnan(v):
+                minval = v
+        return minval
+
+    return nanmin_impl
 
 @overload(np.nanmax)
 def np_nanmax(a):
-    dt = determine_dtype(a)
-    if np.issubdtype(dt, np.complexfloating):
-        return complex_nanmax
-    else:
-        return real_nanmax
+    if not isinstance(a, types.Array):
+        return
+    isnan = get_isnan(a.dtype)
+
+    def nanmax_impl(a):
+        if a.size == 0:
+            raise ValueError("nanmin(): empty array")
+        for view in np.nditer(a):
+            maxval = view.item()
+            break
+        for view in np.nditer(a):
+            v = view.item()
+            if not maxval > v and not isnan(v):
+                maxval = v
+        return maxval
+
+    return nanmax_impl
 
 if numpy_version >= (1, 8):
     @overload(np.nanmean)
@@ -804,6 +725,14 @@ def prepare_ptp_input(a):
     else:
         return arr
 
+@register_jitable
+def less_than(a, b):
+    return a < b
+
+@register_jitable
+def greater_than(a, b):
+    return a > b
+
 def _compute_current_val_impl_gen(op):
     def _compute_current_val_impl(current_val, val):
         if isinstance(current_val, types.Complex):
@@ -1715,7 +1644,7 @@ def determine_dtype(array_like):
     array_like_dt = np.float64
     if isinstance(array_like, types.Array):
         array_like_dt = as_dtype(array_like.dtype)
-    elif isinstance(array_like, (types.Number, types.Boolean)):
+    elif isinstance(array_like, types.Number):
         array_like_dt = as_dtype(array_like)
     elif isinstance(array_like, (types.UniTuple, types.Tuple)):
         coltypes = set()
@@ -2227,69 +2156,11 @@ def array_where(context, builder, sig, args):
     return impl_ret_untracked(context, builder, sig.return_type, res)
 
 
-@register_jitable
-def _where_x_y_scalar(cond, x, y, res):
-    for idx, c in np.ndenumerate(cond):
-        res[idx] = x if c else y
-    return res
-
-
-@register_jitable
-def _where_x_scalar(cond, x, y, res):
-    for idx, c in np.ndenumerate(cond):
-        res[idx] = x if c else y[idx]
-    return res
-
-
-@register_jitable
-def _where_y_scalar(cond, x, y, res):
-    for idx, c in np.ndenumerate(cond):
-        res[idx] = x[idx] if c else y
-    return res
-
-
-def _where_inner(context, builder, sig, args, impl):
-    cond, x, y = sig.args
-
-    x_dt = determine_dtype(x)
-    y_dt = determine_dtype(y)
-    npty = np.promote_types(x_dt, y_dt)
-
-    if cond.layout == 'F':
-        def where_impl(cond, x, y):
-            res = np.asfortranarray(np.empty(cond.shape, dtype=npty))
-            return impl(cond, x, y, res)
-    else:
-        def where_impl(cond, x, y):
-            res = np.empty(cond.shape, dtype=npty)
-            return impl(cond, x, y, res)
-
-    res = context.compile_internal(builder, where_impl, sig, args)
-    return impl_ret_untracked(context, builder, sig.return_type, res)
-
-
-array_scalar_scalar_where = partial(_where_inner, impl=_where_x_y_scalar)
-array_array_scalar_where = partial(_where_inner, impl=_where_y_scalar)
-array_scalar_array_where = partial(_where_inner, impl=_where_x_scalar)
-
-
 @lower_builtin(np.where, types.Any, types.Any, types.Any)
 def any_where(context, builder, sig, args):
-    cond, x, y = sig.args
-
+    cond = sig.args[0]
     if isinstance(cond, types.Array):
-        if isinstance(x, types.Array):
-            if isinstance(y, types.Array):
-                impl = array_where
-            elif isinstance(y, (types.Number, types.Boolean)):
-                impl = array_array_scalar_where
-        elif isinstance(x, (types.Number, types.Boolean)):
-            if isinstance(y, types.Array):
-                impl = array_scalar_array_where
-            elif isinstance(y, (types.Number, types.Boolean)):
-                impl = array_scalar_scalar_where
-
-        return impl(context, builder, sig, args)
+        return array_where(context, builder, sig, args)
 
     def scalar_where_impl(cond, x, y):
         """
@@ -2444,7 +2315,7 @@ def _searchsorted(func):
         return lo
     return searchsorted_inner
 
-_lt = less_than
+_lt = register_jitable(lambda x, y: x < y)
 _le = register_jitable(lambda x, y: x <= y)
 _searchsorted_left = register_jitable(_searchsorted(_lt))
 _searchsorted_right = register_jitable(_searchsorted(_le))
diff --git a/numba/targets/callconv.py b/numba/targets/callconv.py
index 2bb10e655..48186d7c2 100644
--- a/numba/targets/callconv.py
+++ b/numba/targets/callconv.py
@@ -148,6 +148,13 @@ class BaseCallConv(object):
         arginfo = self._get_arg_packer(argtypes)
         return arginfo.from_arguments(builder, raw_args)
 
+    def _fix_argtypes(self, argtypes):
+        """
+        Fix argument types, removing any omitted arguments.
+        """
+        return tuple(ty for ty in argtypes
+                     if not isinstance(ty, types.Omitted))
+
     def _get_arg_packer(self, argtypes):
         """
         Get an argument packer for the given argument types.
diff --git a/numba/typing/npydecl.py b/numba/typing/npydecl.py
index b5bf803e6..56eb8c558 100644
--- a/numba/typing/npydecl.py
+++ b/numba/typing/npydecl.py
@@ -1147,16 +1147,11 @@ class Where(AbstractTemplate):
                         as_dtype(getattr(args[2], 'dtype', args[2]))))
             if isinstance(cond, types.Array):
                 # array where()
-                if isinstance(x, types.Array) and isinstance(y, types.Array):
-                    if (cond.ndim == x.ndim == y.ndim):
-                        if x.layout == y.layout == cond.layout:
-                            retty = types.Array(retdty, x.ndim, x.layout)
-                        else:
-                            retty = types.Array(retdty, x.ndim, 'C')
-                        return signature(retty, *args)
-                else:
-                    # x and y both scalar
-                    retty = types.Array(retdty, cond.ndim, cond.layout)
+                if (cond.ndim == x.ndim == y.ndim):
+                    if x.layout == y.layout == cond.layout:
+                        retty = types.Array(retdty, x.ndim, x.layout)
+                    else:
+                        retty = types.Array(retdty, x.ndim, 'C')
                     return signature(retty, *args)
             else:
                 # scalar where()
diff --git a/numba/typing/templates.py b/numba/typing/templates.py
index 19230e86b..588169e1c 100644
--- a/numba/typing/templates.py
+++ b/numba/typing/templates.py
@@ -86,13 +86,6 @@ class Signature(object):
             return self
         sig = signature(self.return_type, *self.args[1:],
                         recvr=self.args[0])
-
-        # Adjust the python signature
-        params = list(self.pysig.parameters.values())[1:]
-        sig.pysig = utils.pySignature(
-            parameters=params,
-            return_annotation=self.pysig.return_annotation,
-        )
         return sig
 
     def as_function(self):
@@ -575,7 +568,7 @@ class _OverloadAttributeTemplate(AttributeTemplate):
         Get the compiled dispatcher implementing the attribute for
         the given formal signature.
         """
-        cache_key = context, typ, attr, tuple(sig_args), tuple(sig_kws.items())
+        cache_key = context, typ, attr
         try:
             disp = cls._impl_cache[cache_key]
         except KeyError:
diff --git a/numba/unicode.py b/numba/unicode.py
index 0a43fc2e0..2d0b539ea 100644
--- a/numba/unicode.py
+++ b/numba/unicode.py
@@ -529,7 +529,6 @@ def unicode_getitem(s, idx):
 
 
 @overload(operator.add)
-@overload(operator.iadd)
 def unicode_concat(a, b):
     if isinstance(a, types.UnicodeType) and isinstance(b, types.UnicodeType):
         def concat_impl(a, b):
diff --git a/numba/utils.py b/numba/utils.py
index 5e5e4dcbb..c5fae7478 100644
--- a/numba/utils.py
+++ b/numba/utils.py
@@ -78,12 +78,10 @@ else:
 
 try:
     from inspect import signature as pysignature
-    from inspect import Signature as pySignature
     from inspect import Parameter as pyParameter
 except ImportError:
     try:
         from funcsigs import signature as pysignature
-        from funcsigs import Signature as pySignature
         from funcsigs import Parameter as pyParameter
     except ImportError:
         raise ImportError("please install the 'funcsigs' package "


Test patch:
diff --git a/numba/cuda/tests/cudapy/test_cuda_array_interface.py b/numba/cuda/tests/cudapy/test_cuda_array_interface.py
index 4d4218668..90c095201 100644
--- a/numba/cuda/tests/cudapy/test_cuda_array_interface.py
+++ b/numba/cuda/tests/cudapy/test_cuda_array_interface.py
@@ -170,6 +170,39 @@ class TestCudaArrayInterface(CUDATestCase):
             h_arr[1::2]
         )
 
+    def test_negative_strided_issue(self):
+        # issue #3705
+        h_arr = np.random.random(10)
+        c_arr = cuda.to_device(h_arr)
+
+        def base_offset(orig, sliced):
+            return sliced['data'][0] - orig['data'][0]
+
+        h_ai = h_arr.__array_interface__
+        c_ai = c_arr.__cuda_array_interface__
+
+        h_ai_sliced = h_arr[::-1].__array_interface__
+        c_ai_sliced = c_arr[::-1].__cuda_array_interface__
+
+        # Check data offset is correct
+        self.assertEqual(
+            base_offset(h_ai, h_ai_sliced),
+            base_offset(c_ai, c_ai_sliced),
+        )
+        # Check shape and strides are correct
+        self.assertEqual(h_ai_sliced['shape'], c_ai_sliced['shape'])
+        self.assertEqual(h_ai_sliced['strides'], c_ai_sliced['strides'])
+
+    def test_negative_strided_copy_to_host(self):
+        # issue #3705
+        h_arr = np.random.random(10)
+        c_arr = cuda.to_device(h_arr)
+        sliced = c_arr[::-1]
+        with self.assertRaises(NotImplementedError) as raises:
+            sliced.copy_to_host()
+        expected_msg = 'D->H copy not implemented for negative strides'
+        self.assertIn(expected_msg, str(raises.exception))
+
 
 if __name__ == "__main__":
     unittest.main()
diff --git a/numba/tests/test_array_methods.py b/numba/tests/test_array_methods.py
index ad22423da..53433650e 100644
--- a/numba/tests/test_array_methods.py
+++ b/numba/tests/test_array_methods.py
@@ -639,71 +639,6 @@ class TestArrayMethods(MemoryLeakMixin, TestCase):
         for x in (0, 1, True, False, 2.5, 0j):
             check_scal(x)
 
-    def test_np_where_3_broadcast_x_y_scalar(self):
-        pyfunc = np_where_3
-        cfunc = jit(nopython=True)(pyfunc)
-
-        def check_ok(args):
-            expected = pyfunc(*args)
-            got = cfunc(*args)
-            self.assertPreciseEqual(got, expected)
-
-        def a_variations():
-            a = np.linspace(-2, 4, 20)
-            self.random.shuffle(a)
-            yield a
-            yield a.reshape(2, 5, 2)
-            yield a.reshape(4, 5, order='F')
-            yield a.reshape(2, 5, 2)[::-1]
-
-        for a in a_variations():
-            params = (a > 0, 0, 1)
-            check_ok(params)
-
-            params = (a < 0, np.nan, 1 + 4j)
-            check_ok(params)
-
-            params = (a > 1, True, False)
-            check_ok(params)
-
-    def test_np_where_3_broadcast_x_or_y_scalar(self):
-        pyfunc = np_where_3
-        cfunc = jit(nopython=True)(pyfunc)
-
-        def check_ok(args):
-            condition, x, y = args
-
-            expected = pyfunc(condition, x, y)
-            got = cfunc(condition, x, y)
-            self.assertPreciseEqual(got, expected)
-
-            # swap x and y
-            expected = pyfunc(condition, y, x)
-            got = cfunc(condition, y, x)
-            self.assertPreciseEqual(got, expected)
-
-        def array_permutations():
-            x = np.arange(9).reshape(3, 3)
-            yield x
-            yield x * 1.1
-            yield np.asfortranarray(x)
-            yield x[::-1]
-            yield np.linspace(-10, 10, 60).reshape(3, 4, 5) * 1j
-
-        def scalar_permutations():
-            yield 0
-            yield 4.3
-            yield np.nan
-            yield True
-            yield 8 + 4j
-
-        for x in array_permutations():
-            for y in scalar_permutations():
-                x_mean = np.mean(x)
-                condition = x > x_mean
-                params = (condition, x, y)
-                check_ok(params)
-
     def test_item(self):
         pyfunc = array_item
         cfunc = jit(nopython=True)(pyfunc)
diff --git a/numba/tests/test_array_reductions.py b/numba/tests/test_array_reductions.py
index 945bbefde..6cc2c2cc3 100644
--- a/numba/tests/test_array_reductions.py
+++ b/numba/tests/test_array_reductions.py
@@ -757,82 +757,6 @@ class TestArrayReductions(MemoryLeakMixin, TestCase):
         msg = "zero-size array reduction not possible"
         self.assertIn(msg, str(e.exception))
 
-    def test_min_max_complex_basic(self):
-        pyfuncs = array_min_global, array_max_global
-
-        for pyfunc in pyfuncs:
-            cfunc = jit(nopython=True)(pyfunc)
-
-            def check(a):
-                expected = pyfunc(a)
-                got = cfunc(a)
-                self.assertPreciseEqual(expected, got)
-
-            real = np.linspace(-10, 10, 40)
-            real[:4] = real[-1]
-            imag = real * 2
-            a = real - imag * 1j
-            check(a)
-
-            for _ in range(10):
-                self.random.shuffle(real)
-                self.random.shuffle(imag)
-                dtype = self.random.choice([np.complex64, np.complex128])
-                a = real - imag * 1j
-                a[:4] = a[-1]
-                check(a.astype(dtype))
-
-    def test_nanmin_nanmax_complex_basic(self):
-        pyfuncs = array_nanmin, array_nanmax
-
-        for pyfunc in pyfuncs:
-            cfunc = jit(nopython=True)(pyfunc)
-
-            def check(a):
-                expected = pyfunc(a)
-                got = cfunc(a)
-                self.assertPreciseEqual(expected, got)
-
-            real = np.linspace(-10, 10, 40)
-            real[:4] = real[-1]
-            real[5:9] = np.nan
-            imag = real * 2
-            imag[7:12] = np.nan
-            a = real - imag * 1j
-            check(a)
-
-            for _ in range(10):
-                self.random.shuffle(real)
-                self.random.shuffle(imag)
-                a = real - imag * 1j
-                a[:4] = a[-1]
-                check(a)
-
-    def test_nanmin_nanmax_non_array_inputs(self):
-        pyfuncs = array_nanmin, array_nanmax
-
-        def check(a):
-            expected = pyfunc(a)
-            got = cfunc(a)
-            self.assertPreciseEqual(expected, got)
-
-        def a_variations():
-            yield [1, 6, 4, 2]
-            yield ((-10, 4, -12), (5, 200, -30))
-            yield np.array(3)
-            yield (2,)
-            yield 3.142
-            yield False
-            yield (np.nan, 3.142, -5.2, 3.0)
-            yield [np.inf, np.nan, -np.inf]
-            yield [(np.nan, 1.1), (-4.4, 8.7)]
-
-        for pyfunc in pyfuncs:
-            cfunc = jit(nopython=True)(pyfunc)
-
-            for a in a_variations():
-                check(a)
-
     @classmethod
     def install_generated_tests(cls):
         # These form a testing product where each of the combinations are tested
@@ -845,16 +769,16 @@ class TestArrayReductions(MemoryLeakMixin, TestCase):
                            array_std, array_std_global,
                            array_all, array_all_global,
                            array_any, array_any_global,
-                           array_min, array_min_global,
-                           array_max, array_max_global,
-                           array_nanmax, array_nanmin,
                            array_nansum,
                            ]
 
         # these functions only work in real space as no complex comparison
         # operator is implemented
-        reduction_funcs_rspace = [array_argmin, array_argmin_global,
-                                  array_argmax, array_argmax_global]
+        reduction_funcs_rspace = [array_min, array_min_global,
+                                  array_max, array_max_global,
+                                  array_argmin, array_argmin_global,
+                                  array_argmax, array_argmax_global,
+                                  array_nanmax, array_nanmin]
 
         if np_version >= (1, 8):
             reduction_funcs += [array_nanmean, array_nanstd, array_nanvar]
diff --git a/numba/tests/test_dummyarray.py b/numba/tests/test_dummyarray.py
index dfc8f00ad..a964ea4dc 100644
--- a/numba/tests/test_dummyarray.py
+++ b/numba/tests/test_dummyarray.py
@@ -117,6 +117,44 @@ class TestSlicing(unittest.TestCase):
             self.assertEqual(got.shape, expect.shape)
             self.assertEqual(got.strides, expect.strides)
 
+    #### Strided
+
+    def test_strided_1d(self):
+        nparr = np.empty(4)
+        arr = Array.from_desc(0, nparr.shape, nparr.strides,
+                              nparr.dtype.itemsize)
+        xx = -2, -1, 1, 2
+        for x in xx:
+            expect = nparr[::x]
+            got = arr[::x]
+            self.assertSameContig(got, expect)
+            self.assertEqual(got.shape, expect.shape)
+            self.assertEqual(got.strides, expect.strides)
+
+    def test_strided_2d(self):
+        nparr = np.empty((4, 5))
+        arr = Array.from_desc(0, nparr.shape, nparr.strides,
+                              nparr.dtype.itemsize)
+        xx = -2, -1, 1, 2
+        for a, b in itertools.product(xx, xx):
+            expect = nparr[::a, ::b]
+            got = arr[::a, ::b]
+            self.assertSameContig(got, expect)
+            self.assertEqual(got.shape, expect.shape)
+            self.assertEqual(got.strides, expect.strides)
+
+    def test_strided_3d(self):
+        nparr = np.empty((4, 5, 6))
+        arr = Array.from_desc(0, nparr.shape, nparr.strides,
+                              nparr.dtype.itemsize)
+        xx = -2, -1, 1, 2
+        for a, b, c in itertools.product(xx, xx, xx):
+            expect = nparr[::a, ::b, ::c]
+            got = arr[::a, ::b, ::c]
+            self.assertSameContig(got, expect)
+            self.assertEqual(got.shape, expect.shape)
+            self.assertEqual(got.strides, expect.strides)
+
 
 class TestReshape(unittest.TestCase):
     def test_reshape_2d2d(self):
@@ -190,7 +228,7 @@ class TestSqueeze(unittest.TestCase):
             arr.squeeze(axis=1)
         with self.assertRaises(ValueError):
             arr.squeeze(axis=(2, 3))
-        
+
 
 class TestExtent(unittest.TestCase):
     def test_extent_1d(self):
diff --git a/numba/tests/test_extending.py b/numba/tests/test_extending.py
index 8c523af77..cddf1770e 100644
--- a/numba/tests/test_extending.py
+++ b/numba/tests/test_extending.py
@@ -12,7 +12,7 @@ import re
 import numpy as np
 
 from numba import unittest_support as unittest
-from numba import njit, jit, types, errors, typing, compiler
+from numba import jit, types, errors, typing, compiler
 from numba.targets.registry import cpu_target
 from numba.compiler import compile_isolated
 from .support import (TestCase, captured_stdout, tag, temp_directory,
@@ -867,43 +867,6 @@ class TestHighLevelExtending(TestCase):
         self.assertIn("use of VAR_KEYWORD (e.g. **kwargs) is unsupported", msg)
         self.assertIn("offending argument name is '**kws'", msg)
 
-    def test_overload_method_kwargs(self):
-        # Issue #3489
-        @overload_method(types.Array, 'foo')
-        def fooimpl(arr, a_kwarg=10):
-            def impl(arr, a_kwarg=10):
-                return a_kwarg
-            return impl
-
-        @njit
-        def bar(A):
-            return A.foo(), A.foo(20), A.foo(a_kwarg=30)
-
-        Z = np.arange(5)
-
-        self.assertEqual(bar(Z), (10, 20, 30))
-
-    def test_overload_method_literal_unpack(self):
-        # Issue #3683
-        @overload_method(types.Array, 'litfoo')
-        def litfoo(arr, val):
-            # Must be an integer
-            if isinstance(val, types.Integer):
-                # Must not be literal
-                if not isinstance(val, types.Literal):
-                    def impl(arr, val):
-                        return val
-                    return impl
-
-        @njit
-        def bar(A):
-            return A.litfoo(0xcafe)
-
-        A = np.zeros(1)
-        bar(A)
-        self.assertEqual(bar(A), 0xcafe)
-
-
 def _assert_cache_stats(cfunc, expect_hit, expect_misses):
     hit = cfunc._cache_hits[cfunc.signatures[0]]
     if hit != expect_hit:
diff --git a/numba/tests/test_parfors.py b/numba/tests/test_parfors.py
index ffef5539c..29299994e 100644
--- a/numba/tests/test_parfors.py
+++ b/numba/tests/test_parfors.py
@@ -1378,18 +1378,6 @@ class TestParfors(TestParforsBase):
 
         self.check(test_impl, np.random.ranf(128))
 
-    @skip_unsupported
-    def test_array_compare_scalar(self):
-        """ issue3671: X != 0 becomes an arrayexpr with operator.ne.
-            That is turned into a parfor by devectorizing.  Make sure
-            the return type of the devectorized operator.ne
-            on integer types works properly.
-        """
-        def test_impl():
-            X = np.zeros(10, dtype=np.int_)
-            return X != 0
-
-        self.check(test_impl)
 
 class TestPrangeBase(TestParforsBase):
 
@@ -2039,28 +2027,6 @@ class TestPrange(TestPrangeBase):
                         "in non-deterministic or unintended results.")
         self.assertIn(expected_msg, str(warning_obj.message))
 
-    @skip_unsupported
-    def test_nested_parfor_push_call_vars(self):
-        """ issue 3686: if a prange has something inside it that causes
-            a nested parfor to be generated and both the inner and outer
-            parfor use the same call variable defined outside the parfors
-            then ensure that when that call variable is pushed into the
-            parfor that the call variable isn't duplicated with the same
-            name resulting in a redundant type lock.
-        """
-        def test_impl():
-            B = 0
-            f = np.negative
-            for i in range(1):
-                this_matters = f(1.)
-                B += f(np.zeros(1,))[0]
-            for i in range(2):
-                this_matters = f(1.)
-                B += f(np.zeros(1,))[0]
-
-            return B
-        self.prange_tester(test_impl)
-
     @skip_unsupported
     def test_multiple_call_getattr_object(self):
         def test_impl(n):
diff --git a/numba/tests/test_unicode.py b/numba/tests/test_unicode.py
index 7bd859179..83d79fc0a 100644
--- a/numba/tests/test_unicode.py
+++ b/numba/tests/test_unicode.py
@@ -38,9 +38,6 @@ def getitem_usecase(x, i):
 def concat_usecase(x, y):
     return x + y
 
-def inplace_concat_usecase(x, y):
-    x += y
-    return x
 
 def in_usecase(x, y):
     return x in y
@@ -294,15 +291,6 @@ class TestUnicode(BaseTest):
                                  cfunc(a, b),
                                  "'%s' + '%s'?" % (a, b))
 
-    def test_inplace_concat(self, flags=no_pyobj_flags):
-        pyfunc = inplace_concat_usecase
-        cfunc = njit(pyfunc)
-        for a in UNICODE_EXAMPLES:
-            for b in UNICODE_EXAMPLES[::-1]:
-                self.assertEqual(pyfunc(a, b),
-                                 cfunc(a, b),
-                                 "'%s' + '%s'?" % (a, b))
-
     def test_pointless_slice(self, flags=no_pyobj_flags):
         def pyfunc(a):
             return a[:]