From 4e5653e484e0fb028624caa54f61a8d6d84f0357 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 9 Aug 2019 14:17:07 -0600 Subject: [PATCH 01/33] Add the context argument to get_item_pointer and get_item_pointer2 This is required to be able to raise exceptions, for instance, for bounds checking. --- numba/cgutils.py | 4 ++-- numba/cuda/cudaimpl.py | 4 ++-- numba/roc/hsaimpl.py | 2 +- numba/targets/arraymath.py | 6 +++--- numba/targets/arrayobj.py | 34 +++++++++++++++++----------------- numba/targets/npyimpl.py | 3 ++- 6 files changed, 27 insertions(+), 26 deletions(-) diff --git a/numba/cgutils.py b/numba/cgutils.py index ce885159f28..ce8b4fa7b07 100644 --- a/numba/cgutils.py +++ b/numba/cgutils.py @@ -656,10 +656,10 @@ def unpack_tuple(builder, tup, count=None): return vals -def get_item_pointer(builder, aryty, ary, inds, wraparound=False): +def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False): shapes = unpack_tuple(builder, ary.shape, count=aryty.ndim) strides = unpack_tuple(builder, ary.strides, count=aryty.ndim) - return get_item_pointer2(builder, data=ary.data, shape=shapes, + return get_item_pointer2(context, builder, data=ary.data, shape=shapes, strides=strides, layout=aryty.layout, inds=inds, wraparound=wraparound) diff --git a/numba/cuda/cudaimpl.py b/numba/cuda/cudaimpl.py index 7eaf191f3bc..4c0c8ca4cf2 100644 --- a/numba/cuda/cudaimpl.py +++ b/numba/cuda/cudaimpl.py @@ -514,7 +514,7 @@ def imp(context, builder, sig, args): (aryty.ndim, len(indty))) lary = context.make_array(aryty)(context, builder, ary) - ptr = cgutils.get_item_pointer(builder, aryty, lary, indices) + ptr = cgutils.get_item_pointer(context, builder, aryty, lary, indices) # dispatcher to implementation base on dtype return dispatch_fn(context, builder, dtype, ptr, val) return imp @@ -579,7 +579,7 @@ def ptx_atomic_cas_tuple(context, builder, sig, args): lary = context.make_array(aryty)(context, builder, ary) zero = context.get_constant(types.intp, 0) - ptr = cgutils.get_item_pointer(builder, aryty, lary, (zero,)) + ptr = cgutils.get_item_pointer(context, builder, aryty, lary, (zero,)) if aryty.dtype == types.int32: lmod = builder.module return builder.call(nvvmutils.declare_atomic_cas_int32(lmod), (ptr, old, val)) diff --git a/numba/roc/hsaimpl.py b/numba/roc/hsaimpl.py index 961ac0522c0..0e97a7b6188 100644 --- a/numba/roc/hsaimpl.py +++ b/numba/roc/hsaimpl.py @@ -232,7 +232,7 @@ def hsail_atomic_add_tuple(context, builder, sig, args): (aryty.ndim, len(indty))) lary = context.make_array(aryty)(context, builder, ary) - ptr = cgutils.get_item_pointer(builder, aryty, lary, indices) + ptr = cgutils.get_item_pointer(context, builder, aryty, lary, indices) return builder.atomic_rmw("add", ptr, val, ordering='monotonic') diff --git a/numba/targets/arraymath.py b/numba/targets/arraymath.py index 5135254dda1..4c17933fdd6 100644 --- a/numba/targets/arraymath.py +++ b/numba/targets/arraymath.py @@ -2834,7 +2834,7 @@ def array_nonzero(context, builder, sig, args): one = context.get_constant(types.intp, 1) count = cgutils.alloca_once_value(builder, zero) with cgutils.loop_nest(builder, shape, zero.type) as indices: - ptr = cgutils.get_item_pointer2(builder, data, shape, strides, + ptr = cgutils.get_item_pointer2(context, builder, data, shape, strides, layout, indices) val = load_item(context, builder, aryty, ptr) nz = context.is_true(builder, aryty.dtype, val) @@ -2851,7 +2851,7 @@ def array_nonzero(context, builder, sig, args): # And fill them up index = cgutils.alloca_once_value(builder, zero) with cgutils.loop_nest(builder, shape, zero.type) as indices: - ptr = cgutils.get_item_pointer2(builder, data, shape, strides, + ptr = cgutils.get_item_pointer2(context, builder, data, shape, strides, layout, indices) val = load_item(context, builder, aryty, ptr) nz = context.is_true(builder, aryty.dtype, val) @@ -2862,7 +2862,7 @@ def array_nonzero(context, builder, sig, args): indices = (zero,) cur = builder.load(index) for i in range(nouts): - ptr = cgutils.get_item_pointer2(builder, out_datas[i], + ptr = cgutils.get_item_pointer2(context, builder, out_datas[i], out_shape, (), 'C', [cur]) store_item(context, builder, outaryty, indices[i], ptr) diff --git a/numba/targets/arrayobj.py b/numba/targets/arrayobj.py index dfdce139607..e1ac7984a29 100644 --- a/numba/targets/arrayobj.py +++ b/numba/targets/arrayobj.py @@ -275,7 +275,7 @@ def _getitem_array1d(context, builder, arrayty, array, idx, wraparound): """ Look up and return an element from a 1D array. """ - ptr = cgutils.get_item_pointer(builder, arrayty, array, [idx], + ptr = cgutils.get_item_pointer(context, builder, arrayty, array, [idx], wraparound=wraparound) return load_item(context, builder, arrayty, ptr) @@ -364,7 +364,7 @@ def basic_indexing(context, builder, aryty, ary, index_types, indices): # No need to check wraparound, as negative indices were already # fixed in the loop above. - dataptr = cgutils.get_item_pointer(builder, aryty, ary, + dataptr = cgutils.get_item_pointer(context, builder, aryty, ary, output_indices, wraparound=False) return (dataptr, output_shapes, output_strides) @@ -1015,7 +1015,7 @@ def fancy_getitem(context, builder, sig, args, # No need to check for wraparound, as the indexers all ensure # a positive index is returned. - ptr = cgutils.get_item_pointer2(builder, data, shapes, strides, + ptr = cgutils.get_item_pointer(context, builder, data, shapes, strides, aryty.layout, indices, wraparound=False) val = load_item(context, builder, aryty, ptr) @@ -1150,10 +1150,10 @@ def maybe_copy_source(context, builder, use_copy, intp_t = context.get_value_type(types.intp) with cgutils.loop_nest(builder, src_shapes, intp_t) as indices: - src_ptr = cgutils.get_item_pointer2(builder, src_data, + src_ptr = cgutils.get_item_pointer(context, builder, src_data, src_shapes, src_strides, srcty.layout, indices) - dest_ptr = cgutils.get_item_pointer2(builder, data, + dest_ptr = cgutils.get_item_pointer(context, builder, data, copy_shapes, copy_strides, copy_layout, indices) builder.store(builder.load(src_ptr), dest_ptr) @@ -1164,14 +1164,14 @@ def src_getitem(source_indices): with builder.if_else(use_copy, likely=False) as (if_copy, otherwise): with if_copy: builder.store( - cgutils.get_item_pointer2(builder, builder.load(copy_data), + cgutils.get_item_pointer(context, builder, builder.load(copy_data), copy_shapes, copy_strides, copy_layout, source_indices, wraparound=False), src_ptr) with otherwise: builder.store( - cgutils.get_item_pointer2(builder, src_data, + cgutils.get_item_pointer(context, builder, src_data, src_shapes, src_strides, srcty.layout, source_indices, wraparound=False), @@ -1374,7 +1374,7 @@ def src_cleanup(): # No need to check for wraparound, as the indexers all ensure # a positive index is returned. - dest_ptr = cgutils.get_item_pointer2(builder, dest_data, + dest_ptr = cgutils.get_item_pointer(context, builder, dest_data, dest_shapes, dest_strides, aryty.layout, dest_indices, wraparound=False) @@ -2566,7 +2566,7 @@ class IndexedSubIter(BaseSubIter): def compute_pointer(self, context, builder, indices, arrty, arr): assert len(indices) == self.ndim - return cgutils.get_item_pointer(builder, arrty, arr, + return cgutils.get_item_pointer(context, builder, arrty, arr, indices, wraparound=False) class ZeroDimSubIter(BaseSubIter): @@ -3044,7 +3044,7 @@ def _ptr_for_index(self, context, builder, arrty, arr, index): index = builder.udiv(index, shapes[dim]) indices.reverse() - ptr = cgutils.get_item_pointer2(builder, arr.data, shapes, + ptr = cgutils.get_item_pointer(context, builder, arr.data, shapes, strides, arrty.layout, indices) return ptr @@ -3804,10 +3804,10 @@ def _array_copy(context, builder, sig, args): intp_t = context.get_value_type(types.intp) with cgutils.loop_nest(builder, shapes, intp_t) as indices: - src_ptr = cgutils.get_item_pointer2(builder, src_data, + src_ptr = cgutils.get_item_pointer(context, builder, src_data, shapes, src_strides, arytype.layout, indices) - dest_ptr = cgutils.get_item_pointer2(builder, dest_data, + dest_ptr = cgutils.get_item_pointer(context, builder, dest_data, shapes, dest_strides, rettype.layout, indices) builder.store(builder.load(src_ptr), dest_ptr) @@ -3913,10 +3913,10 @@ def array_astype(context, builder, sig, args): intp_t = context.get_value_type(types.intp) with cgutils.loop_nest(builder, shapes, intp_t) as indices: - src_ptr = cgutils.get_item_pointer2(builder, src_data, + src_ptr = cgutils.get_item_pointer(context, builder, src_data, shapes, src_strides, arytype.layout, indices) - dest_ptr = cgutils.get_item_pointer2(builder, dest_data, + dest_ptr = cgutils.get_item_pointer(context, builder, dest_data, shapes, dest_strides, rettype.layout, indices) item = load_item(context, builder, arytype, src_ptr) @@ -4127,7 +4127,7 @@ def assign_sequence_to_array(context, builder, data, shapes, strides, """ def assign_item(indices, valty, val): - ptr = cgutils.get_item_pointer2(builder, data, shapes, strides, + ptr = cgutils.get_item_pointer(context, builder, data, shapes, strides, arrty.layout, indices, wraparound=False) val = context.cast(builder, val, valty, arrty.dtype) store_item(context, builder, arrty, val, ptr) @@ -4403,12 +4403,12 @@ def _do_concatenate(context, builder, axis, order=retty.layout) with loop_nest as indices: - src_ptr = cgutils.get_item_pointer2(builder, arr_data, + src_ptr = cgutils.get_item_pointer(context, builder, arr_data, arr_sh, arr_st, arrty.layout, indices) val = load_item(context, builder, arrty, src_ptr) val = context.cast(builder, val, arrty.dtype, retty.dtype) - dest_ptr = cgutils.get_item_pointer2(builder, ret_data, + dest_ptr = cgutils.get_item_pointer(context, builder, ret_data, ret_shapes, ret_strides, retty.layout, indices) store_item(context, builder, retty, val, dest_ptr) diff --git a/numba/targets/npyimpl.py b/numba/targets/npyimpl.py index d31008cdc77..279ffa66fd8 100644 --- a/numba/targets/npyimpl.py +++ b/numba/targets/npyimpl.py @@ -136,7 +136,8 @@ def create_iter_indices(self): return _ArrayIndexingHelper(self, indices) def _load_effective_address(self, indices): - return cgutils.get_item_pointer2(self.builder, + return cgutils.get_item_pointer2(self.context, + self.builder, data=self.data, shape=self.shape, strides=self.strides, From 9dd1ea887c0538a68e2ac059ee4209f942569f5e Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 9 Aug 2019 14:20:25 -0600 Subject: [PATCH 02/33] Start implementing bounds checking So far we need to: - Add tests - Make the error message match object mode - Add the flag to the public API (right now it is enabled by default for testing purposes) --- numba/cgutils.py | 18 ++++++++++++++---- 1 file changed, 14 insertions(+), 4 deletions(-) diff --git a/numba/cgutils.py b/numba/cgutils.py index ce8b4fa7b07..4817f0cdb4a 100644 --- a/numba/cgutils.py +++ b/numba/cgutils.py @@ -656,16 +656,16 @@ def unpack_tuple(builder, tup, count=None): return vals -def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False): +def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False, boundscheck=True): shapes = unpack_tuple(builder, ary.shape, count=aryty.ndim) strides = unpack_tuple(builder, ary.strides, count=aryty.ndim) return get_item_pointer2(context, builder, data=ary.data, shape=shapes, strides=strides, layout=aryty.layout, inds=inds, - wraparound=wraparound) + wraparound=wraparound, boundscheck=boundscheck) -def get_item_pointer2(builder, data, shape, strides, layout, inds, - wraparound=False): +def get_item_pointer2(context, builder, data, shape, strides, layout, inds, + wraparound=False, boundscheck=True): if wraparound: # Wraparound indices = [] @@ -676,6 +676,16 @@ def get_item_pointer2(builder, data, shape, strides, layout, inds, indices.append(selected) else: indices = inds + if boundscheck: + for ind, dimlen in zip(indices, shape): + msg = "index is out of bounds" + out_of_bounds_upper = builder.icmp_signed('>=', ind, dimlen) + with if_unlikely(builder, out_of_bounds_upper): + context.call_conv.return_user_exc(builder, IndexError, (msg,)) + out_of_bounds_lower = builder.icmp_signed('<', ind, ind.type(0)) + with if_unlikely(builder, out_of_bounds_lower): + context.call_conv.return_user_exc(builder, IndexError, (msg,)) + if not indices: # Indexing with empty tuple return builder.gep(data, [int32_t(0)]) From 991e5e1ffb6a0fb79a4472d4570ac94964562722 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 9 Aug 2019 14:35:17 -0600 Subject: [PATCH 03/33] Fix incorrect function name replacement --- numba/targets/arrayobj.py | 28 ++++++++++++++-------------- 1 file changed, 14 insertions(+), 14 deletions(-) diff --git a/numba/targets/arrayobj.py b/numba/targets/arrayobj.py index e1ac7984a29..2ca9fb6ff44 100644 --- a/numba/targets/arrayobj.py +++ b/numba/targets/arrayobj.py @@ -1015,7 +1015,7 @@ def fancy_getitem(context, builder, sig, args, # No need to check for wraparound, as the indexers all ensure # a positive index is returned. - ptr = cgutils.get_item_pointer(context, builder, data, shapes, strides, + ptr = cgutils.get_item_pointer2(context, builder, data, shapes, strides, aryty.layout, indices, wraparound=False) val = load_item(context, builder, aryty, ptr) @@ -1150,10 +1150,10 @@ def maybe_copy_source(context, builder, use_copy, intp_t = context.get_value_type(types.intp) with cgutils.loop_nest(builder, src_shapes, intp_t) as indices: - src_ptr = cgutils.get_item_pointer(context, builder, src_data, + src_ptr = cgutils.get_item_pointer2(context, builder, src_data, src_shapes, src_strides, srcty.layout, indices) - dest_ptr = cgutils.get_item_pointer(context, builder, data, + dest_ptr = cgutils.get_item_pointer2(context, builder, data, copy_shapes, copy_strides, copy_layout, indices) builder.store(builder.load(src_ptr), dest_ptr) @@ -1164,14 +1164,14 @@ def src_getitem(source_indices): with builder.if_else(use_copy, likely=False) as (if_copy, otherwise): with if_copy: builder.store( - cgutils.get_item_pointer(context, builder, builder.load(copy_data), + cgutils.get_item_pointer2(context, builder, builder.load(copy_data), copy_shapes, copy_strides, copy_layout, source_indices, wraparound=False), src_ptr) with otherwise: builder.store( - cgutils.get_item_pointer(context, builder, src_data, + cgutils.get_item_pointer2(context, builder, src_data, src_shapes, src_strides, srcty.layout, source_indices, wraparound=False), @@ -1374,7 +1374,7 @@ def src_cleanup(): # No need to check for wraparound, as the indexers all ensure # a positive index is returned. - dest_ptr = cgutils.get_item_pointer(context, builder, dest_data, + dest_ptr = cgutils.get_item_pointer2(context, builder, dest_data, dest_shapes, dest_strides, aryty.layout, dest_indices, wraparound=False) @@ -3044,7 +3044,7 @@ def _ptr_for_index(self, context, builder, arrty, arr, index): index = builder.udiv(index, shapes[dim]) indices.reverse() - ptr = cgutils.get_item_pointer(context, builder, arr.data, shapes, + ptr = cgutils.get_item_pointer2(context, builder, arr.data, shapes, strides, arrty.layout, indices) return ptr @@ -3804,10 +3804,10 @@ def _array_copy(context, builder, sig, args): intp_t = context.get_value_type(types.intp) with cgutils.loop_nest(builder, shapes, intp_t) as indices: - src_ptr = cgutils.get_item_pointer(context, builder, src_data, + src_ptr = cgutils.get_item_pointer2(context, builder, src_data, shapes, src_strides, arytype.layout, indices) - dest_ptr = cgutils.get_item_pointer(context, builder, dest_data, + dest_ptr = cgutils.get_item_pointer2(context, builder, dest_data, shapes, dest_strides, rettype.layout, indices) builder.store(builder.load(src_ptr), dest_ptr) @@ -3913,10 +3913,10 @@ def array_astype(context, builder, sig, args): intp_t = context.get_value_type(types.intp) with cgutils.loop_nest(builder, shapes, intp_t) as indices: - src_ptr = cgutils.get_item_pointer(context, builder, src_data, + src_ptr = cgutils.get_item_pointer2(context, builder, src_data, shapes, src_strides, arytype.layout, indices) - dest_ptr = cgutils.get_item_pointer(context, builder, dest_data, + dest_ptr = cgutils.get_item_pointer2(context, builder, dest_data, shapes, dest_strides, rettype.layout, indices) item = load_item(context, builder, arytype, src_ptr) @@ -4127,7 +4127,7 @@ def assign_sequence_to_array(context, builder, data, shapes, strides, """ def assign_item(indices, valty, val): - ptr = cgutils.get_item_pointer(context, builder, data, shapes, strides, + ptr = cgutils.get_item_pointer2(context, builder, data, shapes, strides, arrty.layout, indices, wraparound=False) val = context.cast(builder, val, valty, arrty.dtype) store_item(context, builder, arrty, val, ptr) @@ -4403,12 +4403,12 @@ def _do_concatenate(context, builder, axis, order=retty.layout) with loop_nest as indices: - src_ptr = cgutils.get_item_pointer(context, builder, arr_data, + src_ptr = cgutils.get_item_pointer2(context, builder, arr_data, arr_sh, arr_st, arrty.layout, indices) val = load_item(context, builder, arrty, src_ptr) val = context.cast(builder, val, arrty.dtype, retty.dtype) - dest_ptr = cgutils.get_item_pointer(context, builder, ret_data, + dest_ptr = cgutils.get_item_pointer2(context, builder, ret_data, ret_shapes, ret_strides, retty.layout, indices) store_item(context, builder, retty, val, dest_ptr) From 53549776540d7edf09340ff09db568107686ee1d Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Thu, 15 Aug 2019 15:55:14 -0600 Subject: [PATCH 04/33] Fix bounds checking for slices on shape-0 arrays In NumPy with a = np.empty((0, 1)), a[0, 0] raises IndexError, but a[:, 0] does not. It is impossible to distinguish these two cases in get_item_pointer, which both want a pointer to the "start" of the empty array. So I have factored out the bounds checking into a helper function, and called it separately in basic_indexing() (the same as is already done with wraparound correction). That way, we turn of bounds checking for slice indices on shape 0 axes. Note that slice indices are already not bounds checked from above, e.g., a[0:100] works even if a has fewer than 100 elements. It works this way even on Python lists. --- numba/cgutils.py | 20 +++++++++++--------- numba/targets/arrayobj.py | 8 ++++++-- 2 files changed, 17 insertions(+), 11 deletions(-) diff --git a/numba/cgutils.py b/numba/cgutils.py index 4817f0cdb4a..8ed15253852 100644 --- a/numba/cgutils.py +++ b/numba/cgutils.py @@ -656,16 +656,24 @@ def unpack_tuple(builder, tup, count=None): return vals -def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False, boundscheck=True): +def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False, boundscheck=False): shapes = unpack_tuple(builder, ary.shape, count=aryty.ndim) strides = unpack_tuple(builder, ary.strides, count=aryty.ndim) return get_item_pointer2(context, builder, data=ary.data, shape=shapes, strides=strides, layout=aryty.layout, inds=inds, wraparound=wraparound, boundscheck=boundscheck) +def boundscheck(context, builder, ind, dimlen): + msg = "index is out of bounds" + out_of_bounds_upper = builder.icmp_signed('>=', ind, dimlen) + with if_unlikely(builder, out_of_bounds_upper): + context.call_conv.return_user_exc(builder, IndexError, (msg,)) + out_of_bounds_lower = builder.icmp_signed('<', ind, ind.type(0)) + with if_unlikely(builder, out_of_bounds_lower): + context.call_conv.return_user_exc(builder, IndexError, (msg,)) def get_item_pointer2(context, builder, data, shape, strides, layout, inds, - wraparound=False, boundscheck=True): + wraparound=False, boundscheck=False): if wraparound: # Wraparound indices = [] @@ -678,13 +686,7 @@ def get_item_pointer2(context, builder, data, shape, strides, layout, inds, indices = inds if boundscheck: for ind, dimlen in zip(indices, shape): - msg = "index is out of bounds" - out_of_bounds_upper = builder.icmp_signed('>=', ind, dimlen) - with if_unlikely(builder, out_of_bounds_upper): - context.call_conv.return_user_exc(builder, IndexError, (msg,)) - out_of_bounds_lower = builder.icmp_signed('<', ind, ind.type(0)) - with if_unlikely(builder, out_of_bounds_lower): - context.call_conv.return_user_exc(builder, IndexError, (msg,)) + boundscheck(context, builder, ind, dimlen) if not indices: # Indexing with empty tuple diff --git a/numba/targets/arrayobj.py b/numba/targets/arrayobj.py index 2ca9fb6ff44..bf51699e495 100644 --- a/numba/targets/arrayobj.py +++ b/numba/targets/arrayobj.py @@ -311,7 +311,7 @@ def iternext_array(context, builder, sig, args, result): #------------------------------------------------------------------------------- # Basic indexing (with integers and slices only) -def basic_indexing(context, builder, aryty, ary, index_types, indices): +def basic_indexing(context, builder, aryty, ary, index_types, indices, boundscheck=True): """ Perform basic indexing on the given array. A (data pointer, shapes, strides) tuple is returned describing @@ -345,11 +345,15 @@ def basic_indexing(context, builder, aryty, ary, index_types, indices): output_indices.append(slice.start) sh = slicing.get_slice_length(builder, slice) st = slicing.fix_stride(builder, slice, strides[ax]) + nonzero_dim_p = builder.icmp_signed('!=', shapes[ax], shapes[ax].type(0)) + with cgutils.if_likely(builder, nonzero_dim_p): + cgutils.boundscheck(context, builder, slice.start, shapes[ax]) output_shapes.append(sh) output_strides.append(st) elif isinstance(idxty, types.Integer): ind = fix_integer_index(context, builder, idxty, indexval, shapes[ax]) + cgutils.boundscheck(context, builder, ind, shapes[ax]) output_indices.append(ind) else: raise NotImplementedError("unexpected index type: %s" % (idxty,)) @@ -366,7 +370,7 @@ def basic_indexing(context, builder, aryty, ary, index_types, indices): # fixed in the loop above. dataptr = cgutils.get_item_pointer(context, builder, aryty, ary, output_indices, - wraparound=False) + wraparound=False, boundscheck=False) return (dataptr, output_shapes, output_strides) From c5b6d25188461aff431900ec2c966f559e2b7b66 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Thu, 15 Aug 2019 16:06:43 -0600 Subject: [PATCH 05/33] Actually respect the boundscheck flag to basic_indexing() --- numba/targets/arrayobj.py | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/numba/targets/arrayobj.py b/numba/targets/arrayobj.py index bf51699e495..f28cf85d2b5 100644 --- a/numba/targets/arrayobj.py +++ b/numba/targets/arrayobj.py @@ -345,15 +345,17 @@ def basic_indexing(context, builder, aryty, ary, index_types, indices, boundsche output_indices.append(slice.start) sh = slicing.get_slice_length(builder, slice) st = slicing.fix_stride(builder, slice, strides[ax]) - nonzero_dim_p = builder.icmp_signed('!=', shapes[ax], shapes[ax].type(0)) - with cgutils.if_likely(builder, nonzero_dim_p): - cgutils.boundscheck(context, builder, slice.start, shapes[ax]) + if boundscheck: + nonzero_dim_p = builder.icmp_signed('!=', shapes[ax], shapes[ax].type(0)) + with cgutils.if_likely(builder, nonzero_dim_p): + cgutils.boundscheck(context, builder, slice.start, shapes[ax]) output_shapes.append(sh) output_strides.append(st) elif isinstance(idxty, types.Integer): ind = fix_integer_index(context, builder, idxty, indexval, shapes[ax]) - cgutils.boundscheck(context, builder, ind, shapes[ax]) + if boundscheck: + cgutils.boundscheck(context, builder, ind, shapes[ax]) output_indices.append(ind) else: raise NotImplementedError("unexpected index type: %s" % (idxty,)) From 9a802cd07db6a91b8697588747720741814be9a0 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Thu, 15 Aug 2019 18:03:16 -0600 Subject: [PATCH 06/33] Remove boundschecking from slicing Python/NumPy doesn't do any bounds checking on slices. --- numba/targets/arrayobj.py | 4 ---- 1 file changed, 4 deletions(-) diff --git a/numba/targets/arrayobj.py b/numba/targets/arrayobj.py index f28cf85d2b5..676fc795032 100644 --- a/numba/targets/arrayobj.py +++ b/numba/targets/arrayobj.py @@ -345,10 +345,6 @@ def basic_indexing(context, builder, aryty, ary, index_types, indices, boundsche output_indices.append(slice.start) sh = slicing.get_slice_length(builder, slice) st = slicing.fix_stride(builder, slice, strides[ax]) - if boundscheck: - nonzero_dim_p = builder.icmp_signed('!=', shapes[ax], shapes[ax].type(0)) - with cgutils.if_likely(builder, nonzero_dim_p): - cgutils.boundscheck(context, builder, slice.start, shapes[ax]) output_shapes.append(sh) output_strides.append(st) elif isinstance(idxty, types.Integer): From 0b95b7ca8e287dd772b96fce9bff80829ff59efc Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Thu, 15 Aug 2019 19:16:46 -0600 Subject: [PATCH 07/33] Fix the lowered linspace definition to not use an out-of-bounds index when num == 0 --- numba/targets/arrayobj.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/numba/targets/arrayobj.py b/numba/targets/arrayobj.py index 676fc795032..68d9e257d55 100644 --- a/numba/targets/arrayobj.py +++ b/numba/targets/arrayobj.py @@ -3769,6 +3769,8 @@ def numpy_linspace_3(context, builder, sig, args): def linspace(start, stop, num): arr = np.empty(num, dtype) + if num == 0: + return arr div = num - 1 delta = stop - start arr[0] = start From c413f26cc77e9c15476485696ad497348955fa28 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 16 Aug 2019 10:56:45 -0600 Subject: [PATCH 08/33] Add debug printing of the index, shape, and axis for bounds checking This is only temporary, until we can include it in the error message proper. I have enabled it when NUMBA_FULL_TRACEBACKS is set (I don't want to add a new debug flag for it since it should hopefully be removed soon). --- numba/cgutils.py | 20 +++++++++++++++++--- numba/targets/arrayobj.py | 2 +- 2 files changed, 18 insertions(+), 4 deletions(-) diff --git a/numba/cgutils.py b/numba/cgutils.py index 8ed15253852..fad0268c567 100644 --- a/numba/cgutils.py +++ b/numba/cgutils.py @@ -663,13 +663,27 @@ def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False, bound strides=strides, layout=aryty.layout, inds=inds, wraparound=wraparound, boundscheck=boundscheck) -def boundscheck(context, builder, ind, dimlen): +def boundscheck(context, builder, ind, dimlen, axis=None): + # breakpoint() + def _dbg(): + # Remove this when we figure out how to include this information + # in the error message. + if axis is not None: + if isinstance(axis, int): + printf(builder, "debug: IndexError: index %d is out of bounds for axis {} with size %d\n".format(axis), ind, dimlen) + else: + printf(builder, "debug: IndexError: index %d is out of bounds for axis %d with size %d\n".format(axis), ind, axis, dimlen) + else: + printf(builder, "debug: IndexError: index %d is out of bounds for size %d\n", ind, dimlen) + msg = "index is out of bounds" out_of_bounds_upper = builder.icmp_signed('>=', ind, dimlen) with if_unlikely(builder, out_of_bounds_upper): + if config.FULL_TRACEBACKS: _dbg() context.call_conv.return_user_exc(builder, IndexError, (msg,)) out_of_bounds_lower = builder.icmp_signed('<', ind, ind.type(0)) with if_unlikely(builder, out_of_bounds_lower): + if config.FULL_TRACEBACKS: _dbg() context.call_conv.return_user_exc(builder, IndexError, (msg,)) def get_item_pointer2(context, builder, data, shape, strides, layout, inds, @@ -685,8 +699,8 @@ def get_item_pointer2(context, builder, data, shape, strides, layout, inds, else: indices = inds if boundscheck: - for ind, dimlen in zip(indices, shape): - boundscheck(context, builder, ind, dimlen) + for axis, (ind, dimlen) in enumerate(zip(indices, shape)): + boundscheck(context, builder, ind, dimlen, axis) if not indices: # Indexing with empty tuple diff --git a/numba/targets/arrayobj.py b/numba/targets/arrayobj.py index 68d9e257d55..502977ae8b0 100644 --- a/numba/targets/arrayobj.py +++ b/numba/targets/arrayobj.py @@ -351,7 +351,7 @@ def basic_indexing(context, builder, aryty, ary, index_types, indices, boundsche ind = fix_integer_index(context, builder, idxty, indexval, shapes[ax]) if boundscheck: - cgutils.boundscheck(context, builder, ind, shapes[ax]) + cgutils.boundscheck(context, builder, ind, shapes[ax], ax) output_indices.append(ind) else: raise NotImplementedError("unexpected index type: %s" % (idxty,)) From e094f24b4845fe73de7a7b4b271a055b46031de3 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 27 Sep 2019 15:21:53 -0600 Subject: [PATCH 09/33] Remove debug statement --- numba/cgutils.py | 1 - 1 file changed, 1 deletion(-) diff --git a/numba/cgutils.py b/numba/cgutils.py index fad0268c567..12363a89f75 100644 --- a/numba/cgutils.py +++ b/numba/cgutils.py @@ -664,7 +664,6 @@ def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False, bound wraparound=wraparound, boundscheck=boundscheck) def boundscheck(context, builder, ind, dimlen, axis=None): - # breakpoint() def _dbg(): # Remove this when we figure out how to include this information # in the error message. From b0aa9cd1dfd0f40c3729b80ba52b07385f917332 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 27 Sep 2019 16:34:33 -0600 Subject: [PATCH 10/33] Rename the existing (but unused) "boundcheck" option to "boundscheck" --- numba/compiler.py | 6 +++--- numba/cuda/compiler.py | 2 +- numba/targets/base.py | 2 +- numba/targets/cpu.py | 2 +- numba/targets/options.py | 7 +++---- 5 files changed, 9 insertions(+), 10 deletions(-) diff --git a/numba/compiler.py b/numba/compiler.py index 61348829ffb..84b6dcd4e23 100644 --- a/numba/compiler.py +++ b/numba/compiler.py @@ -44,7 +44,7 @@ class Flags(utils.ConfigOptions): 'release_gil': False, 'no_compile': False, 'debuginfo': False, - 'boundcheck': False, + 'boundscheck': False, 'forceinline': False, 'no_cpython_wrapper': False, # Enable automatic parallel optimization, can be fine-tuned by taking @@ -963,8 +963,8 @@ def _make_subtarget(targetctx, flags): subtargetoptions = {} if flags.debuginfo: subtargetoptions['enable_debuginfo'] = True - if flags.boundcheck: - subtargetoptions['enable_boundcheck'] = True + if flags.boundscheck: + subtargetoptions['enable_boundscheck'] = True if flags.nrt: subtargetoptions['enable_nrt'] = True if flags.auto_parallel: diff --git a/numba/cuda/compiler.py b/numba/cuda/compiler.py index 15b88a8464e..e7bf1630f82 100644 --- a/numba/cuda/compiler.py +++ b/numba/cuda/compiler.py @@ -37,7 +37,7 @@ def compile_cuda(pyfunc, return_type, args, debug, inline): flags.set('no_compile') flags.set('no_cpython_wrapper') if debug: - flags.set('boundcheck') + flags.set('boundscheck') flags.set('debuginfo') if inline: flags.set('forceinline') diff --git a/numba/targets/base.py b/numba/targets/base.py index a3dfe278edc..7593a432c51 100644 --- a/numba/targets/base.py +++ b/numba/targets/base.py @@ -197,7 +197,7 @@ class BaseContext(object): DIBuilder = debuginfo.DIBuilder # Bound checking - enable_boundcheck = False + enable_boundscheck = False # NRT enable_nrt = False diff --git a/numba/targets/cpu.py b/numba/targets/cpu.py index 99f354f01cc..56f346d8d4d 100644 --- a/numba/targets/cpu.py +++ b/numba/targets/cpu.py @@ -204,7 +204,7 @@ class CPUTargetOptions(TargetOptions): "nogil": bool, "forceobj": bool, "looplift": bool, - "boundcheck": bool, + "boundscheck": bool, "debug": bool, "_nrt": bool, "no_rewrites": bool, diff --git a/numba/targets/options.py b/numba/targets/options.py index 0acd7a16c54..7f3fe9efb27 100644 --- a/numba/targets/options.py +++ b/numba/targets/options.py @@ -44,15 +44,15 @@ def set_flags(self, flags): if kws.pop('looplift', True): flags.set("enable_looplift") - if kws.pop('boundcheck', False): - flags.set("boundcheck") + if kws.pop('boundscheck', False): + flags.set("boundscheck") if kws.pop('_nrt', True): flags.set("nrt") if kws.pop('debug', config.DEBUGINFO_DEFAULT): flags.set("debuginfo") - flags.set("boundcheck") + flags.set("boundscheck") if kws.pop('nogil', False): flags.set("release_gil") @@ -80,4 +80,3 @@ def set_flags(self, flags): if kws: # Unread options? raise NameError("Unrecognized options: %s" % kws.keys()) - From c9f0f80f16b67bc9f0ffd02d3fc655c4412653a6 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 27 Sep 2019 16:39:36 -0600 Subject: [PATCH 11/33] Respect the boundscheck flag in @jit --- numba/cgutils.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/numba/cgutils.py b/numba/cgutils.py index 12363a89f75..1967ad445bf 100644 --- a/numba/cgutils.py +++ b/numba/cgutils.py @@ -664,6 +664,9 @@ def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False, bound wraparound=wraparound, boundscheck=boundscheck) def boundscheck(context, builder, ind, dimlen, axis=None): + if not context.enable_boundscheck: + return + def _dbg(): # Remove this when we figure out how to include this information # in the error message. From a3521d83e1058876d813eb9888f40ecb8e20c745 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 27 Sep 2019 16:49:17 -0600 Subject: [PATCH 12/33] Add the NUMBA_BOUNDSCHECK environment variable, which overrides the boundscheck flag to @jit --- numba/cgutils.py | 3 ++- numba/config.py | 4 ++++ 2 files changed, 6 insertions(+), 1 deletion(-) diff --git a/numba/cgutils.py b/numba/cgutils.py index 1967ad445bf..e781551c529 100644 --- a/numba/cgutils.py +++ b/numba/cgutils.py @@ -664,7 +664,8 @@ def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False, bound wraparound=wraparound, boundscheck=boundscheck) def boundscheck(context, builder, ind, dimlen, axis=None): - if not context.enable_boundscheck: + if (config.BOUNDSCHECK is None and not context.enable_boundscheck) \ + or config.BOUNDSCHECK == 0: return def _dbg(): diff --git a/numba/config.py b/numba/config.py index bfcb9bfb7e8..fd05d4896fd 100644 --- a/numba/config.py +++ b/numba/config.py @@ -139,6 +139,10 @@ def optional_str(x): # just bold fonts in use. COLOR_SCHEME = _readenv("NUMBA_COLOR_SCHEME", str, "no_color") + # Whether to globally enable bounds checking. The default None means + # to use the value of the flag to @njit. 0 or 1 overrides the flag globally. + BOUNDSCHECK = _readenv("NUMBA_BOUNDSCHECK", int, None) + # Debug flag to control compiler debug print DEBUG = _readenv("NUMBA_DEBUG", int, 0) From 51a7519456789347bd04e273da3bad5da0b52115 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 27 Sep 2019 17:09:24 -0600 Subject: [PATCH 13/33] Add documentation for boundscheck --- docs/source/reference/envvars.rst | 18 +++++++++++++++++- docs/source/reference/jit-compilation.rst | 14 ++++++++++++-- numba/decorators.py | 14 +++++++++++++- 3 files changed, 42 insertions(+), 4 deletions(-) diff --git a/docs/source/reference/envvars.rst b/docs/source/reference/envvars.rst index 4ff38734e26..70785e62fd9 100644 --- a/docs/source/reference/envvars.rst +++ b/docs/source/reference/envvars.rst @@ -35,6 +35,23 @@ for permanent use by adding:: color_scheme: dark_bg +Jit flags +--------- + +These variables globally override flags to the :func:`~numba.jit` decorator. + +.. envvar:: NUMBA_BOUNDSCHECK + + If set to 0 or 1, globally disable or enable bounds checking, respectively. + The default if the variable is not set or set to an empty string is to use + the ``boundscheck`` flag passed to the :func:`~numba.jit` decorator for a + given function. See the documentation of :ref:`@jit + ` for more information. + + Note, due to limitations in numba, the bounds checking currently produces + exception messages that do not match those from NumPy. If you set + ``NUMBA_FULL_TRACEBACKS=1``, the full exception message with the axis, + index, and shape information will be printed to the terminal. Debugging --------- @@ -347,4 +364,3 @@ Threading Control * ``tbb`` - A threading layer backed by Intel TBB. * ``omp`` - A threading layer backed by OpenMP. * ``workqueue`` - A simple built-in work-sharing task scheduler. - diff --git a/docs/source/reference/jit-compilation.rst b/docs/source/reference/jit-compilation.rst index ab17dbe8aba..0557f0e8152 100644 --- a/docs/source/reference/jit-compilation.rst +++ b/docs/source/reference/jit-compilation.rst @@ -7,7 +7,7 @@ JIT functions .. _jit-decorator: -.. decorator:: numba.jit(signature=None, nopython=False, nogil=False, cache=False, forceobj=False, parallel=False, error_model='python', fastmath=False, locals={}) +.. decorator:: numba.jit(signature=None, nopython=False, nogil=False, cache=False, forceobj=False, parallel=False, error_model='python', fastmath=False, locals={}, boundscheck=False) Compile the decorated function on-the-fly to produce efficient machine code. All parameters all optional. @@ -74,7 +74,7 @@ JIT functions .. _jit-decorator-parallel: If true, *parallel* enables the automatic parallelization of a number of - common Numpy constructs as well as the fusion of adjacent parallel + common Numpy constructs as well as the fusion of adjacent parallel operations to maximize cache locality. The *error_model* option controls the divide-by-zero behavior. @@ -95,6 +95,16 @@ JIT functions accurate versions of some math intrinsics are used (answers to within ``4 ULP``). + .. _jit-decorator-boundscheck: + + If True, ``boundscheck`` enables bounds checking for array indices. Out of + bounds accesses will raise IndexError. The default is to not do bounds + checking. If bounds checking is disabled, out of bounds accesses can + produce garbage results or segfaults. However, enabling bounds checking + will slow down typical functions, so it is recommended to only use this + flag for debugging. You can also set the `NUMBA_BOUNDSCHECK` environment + variable to 0 or 1 to globally override this flag. + The *locals* dictionary may be used to force the :ref:`numba-types` of particular local variables, for example if you want to force the use of single precision floats at some point. In general, we recommend diff --git a/numba/decorators.py b/numba/decorators.py index c49cbbca06a..e163215af68 100644 --- a/numba/decorators.py +++ b/numba/decorators.py @@ -35,7 +35,7 @@ def autojit(*args, **kws): "positional argument.") def jit(signature_or_function=None, locals={}, target='cpu', cache=False, - pipeline_class=None, **options): + pipeline_class=None, boundscheck=False, **options): """ This decorator is used to compile a Python function into native code. @@ -94,6 +94,16 @@ def jit(signature_or_function=None, locals={}, target='cpu', cache=False, NOTE: This inlining is performed at the Numba IR level and is in no way related to LLVM inlining. + boundscheck: bool + Set to True to enable bounds checking for array indices. Out + of bounds accesses will raise IndexError. The default is to + not do bounds checking. If bounds checking is disabled, out of + bounds accesses can produce garbage results or segfaults. + However, enabling bounds checking will slow down typical + functions, so it is recommended to only use this flag for + debugging. You can also set the NUMBA_BOUNDSCHECK environment + variable to 0 or 1 to globally override this flag. + Returns -------- A callable usable as a compiled function. Actual compiling will be @@ -144,6 +154,8 @@ def bar(x, y): if 'restype' in options: raise DeprecationError(_msg_deprecated_signature_arg.format('restype')) + options['boundscheck'] = boundscheck + # Handle signature if signature_or_function is None: # No signature, no function From 8638f8cc6939781d98bec40aa77c6092529b75d5 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Tue, 1 Oct 2019 11:53:39 -0600 Subject: [PATCH 14/33] Fix the boundscheck option in the ufunc builder --- numba/npyufunc/ufuncbuilder.py | 1 + 1 file changed, 1 insertion(+) diff --git a/numba/npyufunc/ufuncbuilder.py b/numba/npyufunc/ufuncbuilder.py index 97cb73ebe96..a9390907ce4 100644 --- a/numba/npyufunc/ufuncbuilder.py +++ b/numba/npyufunc/ufuncbuilder.py @@ -26,6 +26,7 @@ class UFuncTargetOptions(TargetOptions): OPTIONS = { "nopython" : bool, "forceobj" : bool, + "boundscheck": bool, "fastmath" : FastMathOptions, } From 035daf4d1ce29fb35bf9fb047d5382c24a634ff5 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 11 Oct 2019 15:58:20 -0600 Subject: [PATCH 15/33] Add a test for boundscheck One of the tests currently fails with a memory leak error, but I'm not sure how to fix it. --- numba/tests/test_boundscheck.py | 72 +++++++++++++++++++++++++++++++++ 1 file changed, 72 insertions(+) create mode 100644 numba/tests/test_boundscheck.py diff --git a/numba/tests/test_boundscheck.py b/numba/tests/test_boundscheck.py new file mode 100644 index 00000000000..88bf19d67d4 --- /dev/null +++ b/numba/tests/test_boundscheck.py @@ -0,0 +1,72 @@ +from __future__ import print_function, division, absolute_import + +import numpy as np + +from numba.compiler import compile_isolated, DEFAULT_FLAGS +from numba import typeof +from numba.types import float64 +from numba import unittest_support as unittest +from .support import MemoryLeakMixin + +BOUNDSCHECK_FLAGS = DEFAULT_FLAGS.copy() +BOUNDSCHECK_FLAGS.set('boundscheck', True) + +def basic_array_access(a): + return a[10] + +def slice_array_access(a): + # The first index (slice) is not bounds checked + return a[10:, 10] + + +class TestBoundsCheck(MemoryLeakMixin, unittest.TestCase): + def test_basic_array_boundscheck(self): + a = np.arange(5) + # Check the numpy behavior to make sure the test is correct + with self.assertRaises(IndexError): + # TODO: When we raise the same error message as numpy, test that + # they are the same + basic_array_access(a) + + at = typeof(a) + c_noboundscheck = compile_isolated(basic_array_access, [at], + flags=DEFAULT_FLAGS) + noboundscheck = c_noboundscheck.entry_point + c_boundscheck = compile_isolated(basic_array_access, [at], + flags=BOUNDSCHECK_FLAGS) + boundscheck = c_boundscheck.entry_point + # Check that the default flag doesn't raise + noboundscheck(a) + with self.assertRaises(IndexError): + boundscheck(a) + + def test_slice_array_boundscheck(self): + a = np.ones((5, 5)) + b = np.ones((5, 20)) + with self.assertRaises(IndexError): + # TODO: When we raise the same error message as numpy, test that + # they are the same + slice_array_access(a) + # Out of bounds on a slice doesn't raise + slice_array_access(b) + + at = typeof(a) + rt = float64[:] + c_noboundscheck = compile_isolated(slice_array_access, [at], + return_type=rt, + flags=DEFAULT_FLAGS) + noboundscheck = c_noboundscheck.entry_point + c_boundscheck = compile_isolated(slice_array_access, [at], + return_type=rt, + flags=BOUNDSCHECK_FLAGS) + boundscheck = c_boundscheck.entry_point + # Check that the default flag doesn't raise + noboundscheck(a) + noboundscheck(b) + with self.assertRaises(IndexError): + boundscheck(a) + # Doesn't raise + boundscheck(b) + +if __name__ == '__main__': + unittest.main() From 450ade39340cc6d7fb10e7ba3b7e433e077205a9 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 15 Nov 2019 15:56:58 -0700 Subject: [PATCH 16/33] Rename boundscheck() to do_boundscheck() That way it doesn't overwrite the flag name. Also add comments clarifying how the boundscheck flag to get_item_pointer/get_item_pointer2 should work. --- numba/cgutils.py | 10 ++++++++-- numba/targets/arrayobj.py | 2 +- 2 files changed, 9 insertions(+), 3 deletions(-) diff --git a/numba/cgutils.py b/numba/cgutils.py index e781551c529..23b62b1ac6c 100644 --- a/numba/cgutils.py +++ b/numba/cgutils.py @@ -657,13 +657,16 @@ def unpack_tuple(builder, tup, count=None): def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False, boundscheck=False): + # Set boundscheck=True for any pointer access that should be + # boundschecked. do_boundscheck() will handle enabling or disabling the + # actual boundschecking based on the user config. shapes = unpack_tuple(builder, ary.shape, count=aryty.ndim) strides = unpack_tuple(builder, ary.strides, count=aryty.ndim) return get_item_pointer2(context, builder, data=ary.data, shape=shapes, strides=strides, layout=aryty.layout, inds=inds, wraparound=wraparound, boundscheck=boundscheck) -def boundscheck(context, builder, ind, dimlen, axis=None): +def do_boundscheck(context, builder, ind, dimlen, axis=None): if (config.BOUNDSCHECK is None and not context.enable_boundscheck) \ or config.BOUNDSCHECK == 0: return @@ -691,6 +694,9 @@ def _dbg(): def get_item_pointer2(context, builder, data, shape, strides, layout, inds, wraparound=False, boundscheck=False): + # Set boundscheck=True for any pointer access that should be + # boundschecked. do_boundscheck() will handle enabling or disabling the + # actual boundschecking based on the user config. if wraparound: # Wraparound indices = [] @@ -703,7 +709,7 @@ def get_item_pointer2(context, builder, data, shape, strides, layout, inds, indices = inds if boundscheck: for axis, (ind, dimlen) in enumerate(zip(indices, shape)): - boundscheck(context, builder, ind, dimlen, axis) + do_boundscheck(context, builder, ind, dimlen, axis) if not indices: # Indexing with empty tuple diff --git a/numba/targets/arrayobj.py b/numba/targets/arrayobj.py index 95f75118eb5..c8f7e9a2716 100644 --- a/numba/targets/arrayobj.py +++ b/numba/targets/arrayobj.py @@ -352,7 +352,7 @@ def basic_indexing(context, builder, aryty, ary, index_types, indices, boundsche ind = fix_integer_index(context, builder, idxty, indexval, shapes[ax]) if boundscheck: - cgutils.boundscheck(context, builder, ind, shapes[ax], ax) + cgutils.do_boundscheck(context, builder, ind, shapes[ax], ax) output_indices.append(ind) else: raise NotImplementedError("unexpected index type: %s" % (idxty,)) From b9a78bd6b9c496ac12b6a83ef05b6daddb20f748 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 15 Nov 2019 16:06:25 -0700 Subject: [PATCH 17/33] Enable bounds checking for fancy indexing --- numba/targets/arrayobj.py | 2 +- numba/tests/test_boundscheck.py | 34 ++++++++++++++++++++++++++++++++- 2 files changed, 34 insertions(+), 2 deletions(-) diff --git a/numba/targets/arrayobj.py b/numba/targets/arrayobj.py index c8f7e9a2716..66f92227ad1 100644 --- a/numba/targets/arrayobj.py +++ b/numba/targets/arrayobj.py @@ -1019,7 +1019,7 @@ def fancy_getitem(context, builder, sig, args, # No need to check for wraparound, as the indexers all ensure # a positive index is returned. ptr = cgutils.get_item_pointer2(context, builder, data, shapes, strides, - aryty.layout, indices, wraparound=False) + aryty.layout, indices, wraparound=False, boundscheck=True) val = load_item(context, builder, aryty, ptr) # Since the destination is C-contiguous, no need for multi-dimensional diff --git a/numba/tests/test_boundscheck.py b/numba/tests/test_boundscheck.py index 88bf19d67d4..ff854eeb1d6 100644 --- a/numba/tests/test_boundscheck.py +++ b/numba/tests/test_boundscheck.py @@ -4,7 +4,7 @@ from numba.compiler import compile_isolated, DEFAULT_FLAGS from numba import typeof -from numba.types import float64 +from numba.types import float64, intp from numba import unittest_support as unittest from .support import MemoryLeakMixin @@ -18,6 +18,9 @@ def slice_array_access(a): # The first index (slice) is not bounds checked return a[10:, 10] +def fancy_array_access(x): + a = np.array([1, 2, 3]) + return x[a] class TestBoundsCheck(MemoryLeakMixin, unittest.TestCase): def test_basic_array_boundscheck(self): @@ -68,5 +71,34 @@ def test_slice_array_boundscheck(self): # Doesn't raise boundscheck(b) + def test_fancy_indexing_boundscheck(self): + a = np.arange(3) + b = np.arange(4) + + # Check the numpy behavior to ensure the test is correct. + with self.assertRaises(IndexError): + # TODO: When we raise the same error message as numpy, test that + # they are the same + fancy_array_access(a) + fancy_array_access(b) + + at = typeof(a) + rt = intp[:] + c_noboundscheck = compile_isolated(fancy_array_access, [at], + return_type=rt, + flags=DEFAULT_FLAGS) + noboundscheck = c_noboundscheck.entry_point + c_boundscheck = compile_isolated(fancy_array_access, [at], + return_type=rt, + flags=BOUNDSCHECK_FLAGS) + boundscheck = c_boundscheck.entry_point + # Check that the default flag doesn't raise + noboundscheck(a) + noboundscheck(b) + with self.assertRaises(IndexError): + boundscheck(a) + # Doesn't raise + boundscheck(b) + if __name__ == '__main__': unittest.main() From 9a108084331053cef1846b25243135315a644dbb Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 15 Nov 2019 16:13:13 -0700 Subject: [PATCH 18/33] Split out the boundscheck tests that raise exceptions They leak memory so have to be tested in a separate test case that doesn't have the memory leak mixin. --- numba/tests/test_boundscheck.py | 74 ++++++++++++++++++++++++++++----- 1 file changed, 64 insertions(+), 10 deletions(-) diff --git a/numba/tests/test_boundscheck.py b/numba/tests/test_boundscheck.py index ff854eeb1d6..068fbfb0661 100644 --- a/numba/tests/test_boundscheck.py +++ b/numba/tests/test_boundscheck.py @@ -22,7 +22,7 @@ def fancy_array_access(x): a = np.array([1, 2, 3]) return x[a] -class TestBoundsCheck(MemoryLeakMixin, unittest.TestCase): +class TestBoundsCheckNoError(MemoryLeakMixin, unittest.TestCase): def test_basic_array_boundscheck(self): a = np.arange(5) # Check the numpy behavior to make sure the test is correct @@ -35,13 +35,9 @@ def test_basic_array_boundscheck(self): c_noboundscheck = compile_isolated(basic_array_access, [at], flags=DEFAULT_FLAGS) noboundscheck = c_noboundscheck.entry_point - c_boundscheck = compile_isolated(basic_array_access, [at], - flags=BOUNDSCHECK_FLAGS) - boundscheck = c_boundscheck.entry_point # Check that the default flag doesn't raise noboundscheck(a) - with self.assertRaises(IndexError): - boundscheck(a) + # boundscheck(a) is tested in TestBoundsCheckError below def test_slice_array_boundscheck(self): a = np.ones((5, 5)) @@ -66,8 +62,8 @@ def test_slice_array_boundscheck(self): # Check that the default flag doesn't raise noboundscheck(a) noboundscheck(b) - with self.assertRaises(IndexError): - boundscheck(a) + # boundscheck(a) is tested in TestBoundsCheckError below + # Doesn't raise boundscheck(b) @@ -95,10 +91,68 @@ def test_fancy_indexing_boundscheck(self): # Check that the default flag doesn't raise noboundscheck(a) noboundscheck(b) - with self.assertRaises(IndexError): - boundscheck(a) + # boundscheck(a) is tested in TestBoundsCheckError below + # Doesn't raise boundscheck(b) +# This is a separate test because the jitted functions that raise exceptions +# have memory leaks. +class TestBoundsCheckError(unittest.TestCase): + def test_basic_array_boundscheck(self): + a = np.arange(5) + # Check the numpy behavior to make sure the test is correct + with self.assertRaises(IndexError): + # TODO: When we raise the same error message as numpy, test that + # they are the same + basic_array_access(a) + + at = typeof(a) + c_boundscheck = compile_isolated(basic_array_access, [at], + flags=BOUNDSCHECK_FLAGS) + boundscheck = c_boundscheck.entry_point + + with self.assertRaises(IndexError): + boundscheck(a) + + def test_slice_array_boundscheck(self): + a = np.ones((5, 5)) + b = np.ones((5, 20)) + with self.assertRaises(IndexError): + # TODO: When we raise the same error message as numpy, test that + # they are the same + slice_array_access(a) + # Out of bounds on a slice doesn't raise + slice_array_access(b) + + at = typeof(a) + rt = float64[:] + c_boundscheck = compile_isolated(slice_array_access, [at], + return_type=rt, + flags=BOUNDSCHECK_FLAGS) + boundscheck = c_boundscheck.entry_point + with self.assertRaises(IndexError): + boundscheck(a) + + def test_fancy_indexing_boundscheck(self): + a = np.arange(3) + b = np.arange(4) + + # Check the numpy behavior to ensure the test is correct. + with self.assertRaises(IndexError): + # TODO: When we raise the same error message as numpy, test that + # they are the same + fancy_array_access(a) + fancy_array_access(b) + + at = typeof(a) + rt = intp[:] + c_boundscheck = compile_isolated(fancy_array_access, [at], + return_type=rt, + flags=BOUNDSCHECK_FLAGS) + boundscheck = c_boundscheck.entry_point + with self.assertRaises(IndexError): + boundscheck(a) + if __name__ == '__main__': unittest.main() From 5dd8759e292aed0837138083368323116e8d357f Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 15 Nov 2019 16:35:15 -0700 Subject: [PATCH 19/33] Make the boundscheck tests work even if NUMBA_BOUNDSCHECK is set --- numba/tests/test_boundscheck.py | 16 +++++++++++++++- 1 file changed, 15 insertions(+), 1 deletion(-) diff --git a/numba/tests/test_boundscheck.py b/numba/tests/test_boundscheck.py index 068fbfb0661..47a36e49fe3 100644 --- a/numba/tests/test_boundscheck.py +++ b/numba/tests/test_boundscheck.py @@ -3,7 +3,7 @@ import numpy as np from numba.compiler import compile_isolated, DEFAULT_FLAGS -from numba import typeof +from numba import typeof, config from numba.types import float64, intp from numba import unittest_support as unittest from .support import MemoryLeakMixin @@ -23,6 +23,10 @@ def fancy_array_access(x): return x[a] class TestBoundsCheckNoError(MemoryLeakMixin, unittest.TestCase): + def setUp(self): + self.old_boundscheck = config.BOUNDSCHECK + config.BOUNDSCHECK = None + def test_basic_array_boundscheck(self): a = np.arange(5) # Check the numpy behavior to make sure the test is correct @@ -96,9 +100,16 @@ def test_fancy_indexing_boundscheck(self): # Doesn't raise boundscheck(b) + def tearDown(self): + config.BOUNDSCHECK = self.old_boundscheck + # This is a separate test because the jitted functions that raise exceptions # have memory leaks. class TestBoundsCheckError(unittest.TestCase): + def setUp(self): + self.old_boundscheck = config.BOUNDSCHECK + config.BOUNDSCHECK = None + def test_basic_array_boundscheck(self): a = np.arange(5) # Check the numpy behavior to make sure the test is correct @@ -154,5 +165,8 @@ def test_fancy_indexing_boundscheck(self): with self.assertRaises(IndexError): boundscheck(a) + def tearDown(self): + config.BOUNDSCHECK = self.old_boundscheck + if __name__ == '__main__': unittest.main() From 847bd0a325043c112810881dd7a2b4b915c3d268 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 22 Nov 2019 15:27:08 -0700 Subject: [PATCH 20/33] Disable boundschecking for CUDA --- numba/cgutils.py | 5 +++++ numba/cuda/decorators.py | 4 +++- numba/tests/test_boundscheck.py | 8 +++++++- 3 files changed, 15 insertions(+), 2 deletions(-) diff --git a/numba/cgutils.py b/numba/cgutils.py index 4e265be6098..6c91d79d07a 100644 --- a/numba/cgutils.py +++ b/numba/cgutils.py @@ -668,6 +668,11 @@ def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False, bound wraparound=wraparound, boundscheck=boundscheck) def do_boundscheck(context, builder, ind, dimlen, axis=None): + # Boundschecking is always disabled for CUDA + from .cuda.target import CUDATargetContext + if isinstance(context, CUDATargetContext): + return + if (config.BOUNDSCHECK is None and not context.enable_boundscheck) \ or config.BOUNDSCHECK == 0: return diff --git a/numba/cuda/decorators.py b/numba/cuda/decorators.py index 56e3ff4e105..d5cecef5bfa 100644 --- a/numba/cuda/decorators.py +++ b/numba/cuda/decorators.py @@ -52,6 +52,9 @@ def jit(func_or_sig=None, argtypes=None, device=False, inline=False, bind=True, if link and config.ENABLE_CUDASIM: raise NotImplementedError('Cannot link PTX in the simulator') + if 'boundscheck' in kws: + raise NotImplementedError("bounds checking is not supported for CUDA") + fastmath = kws.get('fastmath', False) if argtypes is None and not sigutils.is_signature(func_or_sig): if func_or_sig is None: @@ -126,4 +129,3 @@ def convert_types(restype, argtypes): argtypes, restype = sigutils.normalize_signature(restype) return restype, argtypes - diff --git a/numba/tests/test_boundscheck.py b/numba/tests/test_boundscheck.py index 47a36e49fe3..0a96df41981 100644 --- a/numba/tests/test_boundscheck.py +++ b/numba/tests/test_boundscheck.py @@ -3,7 +3,7 @@ import numpy as np from numba.compiler import compile_isolated, DEFAULT_FLAGS -from numba import typeof, config +from numba import typeof, config, cuda from numba.types import float64, intp from numba import unittest_support as unittest from .support import MemoryLeakMixin @@ -100,6 +100,12 @@ def test_fancy_indexing_boundscheck(self): # Doesn't raise boundscheck(b) + def test_no_cuda_boundscheck(self): + with self.assertRaises(NotImplementedError): + @cuda.jit(boundscheck=True) + def func(): + pass + def tearDown(self): config.BOUNDSCHECK = self.old_boundscheck From ce1105b24ce6a3164269056d1e616b76c525d82d Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Fri, 22 Nov 2019 17:21:15 -0700 Subject: [PATCH 21/33] Attempt at getting parfor to ignore boundscheck for slices Parfor converts slices to individual get/setitems, so those should not be boundschecked. This is not so simple, however, because parfor generates code at the numba IR level and the boundschecking is implemented at the LLVM IR level. So this implements it by setting a flag on the node, then globally disabling boundschecking when that node is lowered. This is a bit hacky, but I don't know how else to do it. Unfortunately, this doesn't work. It fixes the test in question (test_parfor_slice19), but then a bunch of other parfor tests fail with '@do_scheduling_unsigned' not found in ..., and I'm not really clear why. --- numba/ir.py | 7 ++++--- numba/lowering.py | 29 +++++++++++++++++++++++------ numba/parfor.py | 5 +++-- 3 files changed, 30 insertions(+), 11 deletions(-) diff --git a/numba/ir.py b/numba/ir.py index 5fefef75cd9..ec3137ce4bb 100644 --- a/numba/ir.py +++ b/numba/ir.py @@ -486,12 +486,12 @@ def getattr(cls, value, attr, loc): return cls(op=op, loc=loc, value=value, attr=attr) @classmethod - def getitem(cls, value, index, loc): + def getitem(cls, value, index, loc, boundscheck=None): assert isinstance(value, Var) assert isinstance(index, Var) assert isinstance(loc, Loc) op = 'getitem' - return cls(op=op, loc=loc, value=value, index=index) + return cls(op=op, loc=loc, value=value, index=index, boundscheck=boundscheck) @classmethod def static_getitem(cls, value, index, index_var, loc): @@ -552,7 +552,7 @@ class SetItem(Stmt): target[index] = value """ - def __init__(self, target, index, value, loc): + def __init__(self, target, index, value, loc, boundscheck=None): assert isinstance(target, Var) assert isinstance(index, Var) assert isinstance(value, Var) @@ -561,6 +561,7 @@ def __init__(self, target, index, value, loc): self.index = index self.value = value self.loc = loc + self.boundscheck = boundscheck def __repr__(self): return '%s[%s] = %s' % (self.target, self.index, self.value) diff --git a/numba/lowering.py b/numba/lowering.py index 3fb48c03f62..c1710e45be3 100644 --- a/numba/lowering.py +++ b/numba/lowering.py @@ -357,7 +357,7 @@ def lower_inst(self, inst): signature = self.fndesc.calltypes[inst] assert signature is not None return self.lower_setitem(inst.target, inst.index, inst.value, - signature) + signature, boundscheck=inst.boundscheck) elif isinstance(inst, ir.StoreMap): signature = self.fndesc.calltypes[inst] @@ -417,7 +417,7 @@ def lower_inst(self, inst): return raise NotImplementedError(type(inst)) - def lower_setitem(self, target_var, index_var, value_var, signature): + def lower_setitem(self, target_var, index_var, value_var, signature, boundscheck=None): target = self.loadvar(target_var.name) value = self.loadvar(value_var.name) index = self.loadvar(index_var.name) @@ -445,7 +445,15 @@ def lower_setitem(self, target_var, index_var, value_var, signature): value = self.context.cast(self.builder, value, valuety, signature.args[2]) - return impl(self.builder, (target, index, value)) + + old_boundscheck = config.BOUNDSCHECK + if boundscheck is not None: + config.BOUNDSCHECK = boundscheck + + try: + return impl(self.builder, (target, index, value)) + finally: + config.BOUNDSCHECK = old_boundscheck def lower_static_raise(self, inst): if inst.exc_class is None: @@ -592,7 +600,7 @@ def try_static_impl(tys, args): res = impl(self.builder, (lhs, rhs)) return cast_result(res) - def lower_getitem(self, resty, expr, value, index, signature): + def lower_getitem(self, resty, expr, value, index, signature, boundscheck=None): baseval = self.loadvar(value.name) indexval = self.loadvar(index.name) # Get implementation of getitem @@ -609,7 +617,16 @@ def lower_getitem(self, resty, expr, value, index, signature): castvals = [self.context.cast(self.builder, av, at, ft) for av, at, ft in zip(argvals, argtyps, signature.args)] - res = impl(self.builder, castvals) + + old_boundscheck = config.BOUNDSCHECK + if boundscheck is not None: + config.BOUNDSCHECK = boundscheck + + try: + res = impl(self.builder, castvals) + finally: + config.BOUNDSCHECK = old_boundscheck + return self.context.cast(self.builder, res, signature.return_type, resty) @@ -1057,7 +1074,7 @@ def lower_expr(self, resty, expr): elif expr.op == "getitem": signature = self.fndesc.calltypes[expr] return self.lower_getitem(resty, expr, expr.value, expr.index, - signature) + signature, boundscheck=expr.boundscheck) elif expr.op == "build_tuple": itemvals = [self.loadvar(i.name) for i in expr.items] diff --git a/numba/parfor.py b/numba/parfor.py index 02ff1fa7e38..47ff30c6bb2 100644 --- a/numba/parfor.py +++ b/numba/parfor.py @@ -2231,13 +2231,14 @@ def _setitem_to_parfor(self, equiv_set, loc, target, index, value, shape=None): if isinstance(value_typ, types.npytypes.Array): value_var = ir.Var(scope, mk_unique_var("$value_var"), loc) self.typemap[value_var.name] = value_typ.dtype - getitem_call = ir.Expr.getitem(value, index_var, loc) + # Disable boundschecking, as this could be part of a slice + getitem_call = ir.Expr.getitem(value, index_var, loc, boundscheck=False) self.calltypes[getitem_call] = signature( value_typ.dtype, value_typ, index_var_typ) true_block.body.append(ir.Assign(getitem_call, value_var, loc)) else: value_var = value - setitem_node = ir.SetItem(target, index_var, value_var, loc) + setitem_node = ir.SetItem(target, index_var, value_var, loc, boundscheck=False) self.calltypes[setitem_node] = signature( types.none, self.typemap[target.name], index_var_typ, el_typ) true_block.body.append(setitem_node) From 5a48607574d482aea6c1977324ef39572cafff04 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Mon, 25 Nov 2019 17:03:30 -0700 Subject: [PATCH 22/33] Fix a test that fails with boundschecking enabled It does a diff of the instructions, but this fails with boundschecking enabled because of the different pickled exceptions. We just manually disable boundschecking for the test. --- numba/tests/test_parfors.py | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/numba/tests/test_parfors.py b/numba/tests/test_parfors.py index 41e2b76727a..9f127729643 100644 --- a/numba/tests/test_parfors.py +++ b/numba/tests/test_parfors.py @@ -20,7 +20,7 @@ import numba from numba import unittest_support as unittest -from .support import TestCase, captured_stdout, MemoryLeakMixin +from .support import TestCase, captured_stdout, MemoryLeakMixin, override_env_config from numba import njit, prange, stencil, inline_closurecall from numba import compiler, typing, errors, typed_passes from numba.targets import cpu @@ -2447,10 +2447,12 @@ def unsigned_variant(): A += i return A - signed_asm = self.get_gufunc_asm(signed_variant, 'signed', - fastmath=True) - unsigned_asm = self.get_gufunc_asm(unsigned_variant, 'unsigned', - fastmath=True) + # Boundschecking breaks the diff check below because of the pickled exception + with override_env_config('NUMBA_BOUNDSCHECK', '0'): + signed_asm = self.get_gufunc_asm(signed_variant, 'signed', + fastmath=True) + unsigned_asm = self.get_gufunc_asm(unsigned_variant, 'unsigned', + fastmath=True) def strip_instrs(asm): acc = [] From f18b79f8ff11696e5139ac6e23049bc46c015cd7 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Mon, 25 Nov 2019 17:09:05 -0700 Subject: [PATCH 23/33] Disable bounds checking on a test of vectorization --- numba/tests/test_parfors.py | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/numba/tests/test_parfors.py b/numba/tests/test_parfors.py index 9f127729643..f85acbec764 100644 --- a/numba/tests/test_parfors.py +++ b/numba/tests/test_parfors.py @@ -2400,11 +2400,13 @@ def will_vectorize(A): arg = np.zeros(10) - novec_asm = self.get_gufunc_asm(will_not_vectorize, 'signed', arg, - fastmath=True) + # Boundschecking breaks vectorization + with override_env_config('NUMBA_BOUNDSCHECK', '0'): + novec_asm = self.get_gufunc_asm(will_not_vectorize, 'signed', arg, + fastmath=True) - vec_asm = self.get_gufunc_asm(will_vectorize, 'unsigned', arg, - fastmath=True) + vec_asm = self.get_gufunc_asm(will_vectorize, 'unsigned', arg, + fastmath=True) for v in novec_asm.values(): # vector variant should not be present From 663116af11b31ee1312e76f40bd1efcabf60b969 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Tue, 26 Nov 2019 09:53:26 -0700 Subject: [PATCH 24/33] Revert "Attempt at getting parfor to ignore boundscheck for slices" This reverts commit ce1105b24ce6a3164269056d1e616b76c525d82d. --- numba/ir.py | 7 +++---- numba/lowering.py | 29 ++++++----------------------- numba/parfor.py | 5 ++--- 3 files changed, 11 insertions(+), 30 deletions(-) diff --git a/numba/ir.py b/numba/ir.py index ec3137ce4bb..5fefef75cd9 100644 --- a/numba/ir.py +++ b/numba/ir.py @@ -486,12 +486,12 @@ def getattr(cls, value, attr, loc): return cls(op=op, loc=loc, value=value, attr=attr) @classmethod - def getitem(cls, value, index, loc, boundscheck=None): + def getitem(cls, value, index, loc): assert isinstance(value, Var) assert isinstance(index, Var) assert isinstance(loc, Loc) op = 'getitem' - return cls(op=op, loc=loc, value=value, index=index, boundscheck=boundscheck) + return cls(op=op, loc=loc, value=value, index=index) @classmethod def static_getitem(cls, value, index, index_var, loc): @@ -552,7 +552,7 @@ class SetItem(Stmt): target[index] = value """ - def __init__(self, target, index, value, loc, boundscheck=None): + def __init__(self, target, index, value, loc): assert isinstance(target, Var) assert isinstance(index, Var) assert isinstance(value, Var) @@ -561,7 +561,6 @@ def __init__(self, target, index, value, loc, boundscheck=None): self.index = index self.value = value self.loc = loc - self.boundscheck = boundscheck def __repr__(self): return '%s[%s] = %s' % (self.target, self.index, self.value) diff --git a/numba/lowering.py b/numba/lowering.py index c1710e45be3..3fb48c03f62 100644 --- a/numba/lowering.py +++ b/numba/lowering.py @@ -357,7 +357,7 @@ def lower_inst(self, inst): signature = self.fndesc.calltypes[inst] assert signature is not None return self.lower_setitem(inst.target, inst.index, inst.value, - signature, boundscheck=inst.boundscheck) + signature) elif isinstance(inst, ir.StoreMap): signature = self.fndesc.calltypes[inst] @@ -417,7 +417,7 @@ def lower_inst(self, inst): return raise NotImplementedError(type(inst)) - def lower_setitem(self, target_var, index_var, value_var, signature, boundscheck=None): + def lower_setitem(self, target_var, index_var, value_var, signature): target = self.loadvar(target_var.name) value = self.loadvar(value_var.name) index = self.loadvar(index_var.name) @@ -445,15 +445,7 @@ def lower_setitem(self, target_var, index_var, value_var, signature, boundscheck value = self.context.cast(self.builder, value, valuety, signature.args[2]) - - old_boundscheck = config.BOUNDSCHECK - if boundscheck is not None: - config.BOUNDSCHECK = boundscheck - - try: - return impl(self.builder, (target, index, value)) - finally: - config.BOUNDSCHECK = old_boundscheck + return impl(self.builder, (target, index, value)) def lower_static_raise(self, inst): if inst.exc_class is None: @@ -600,7 +592,7 @@ def try_static_impl(tys, args): res = impl(self.builder, (lhs, rhs)) return cast_result(res) - def lower_getitem(self, resty, expr, value, index, signature, boundscheck=None): + def lower_getitem(self, resty, expr, value, index, signature): baseval = self.loadvar(value.name) indexval = self.loadvar(index.name) # Get implementation of getitem @@ -617,16 +609,7 @@ def lower_getitem(self, resty, expr, value, index, signature, boundscheck=None): castvals = [self.context.cast(self.builder, av, at, ft) for av, at, ft in zip(argvals, argtyps, signature.args)] - - old_boundscheck = config.BOUNDSCHECK - if boundscheck is not None: - config.BOUNDSCHECK = boundscheck - - try: - res = impl(self.builder, castvals) - finally: - config.BOUNDSCHECK = old_boundscheck - + res = impl(self.builder, castvals) return self.context.cast(self.builder, res, signature.return_type, resty) @@ -1074,7 +1057,7 @@ def lower_expr(self, resty, expr): elif expr.op == "getitem": signature = self.fndesc.calltypes[expr] return self.lower_getitem(resty, expr, expr.value, expr.index, - signature, boundscheck=expr.boundscheck) + signature) elif expr.op == "build_tuple": itemvals = [self.loadvar(i.name) for i in expr.items] diff --git a/numba/parfor.py b/numba/parfor.py index 47ff30c6bb2..02ff1fa7e38 100644 --- a/numba/parfor.py +++ b/numba/parfor.py @@ -2231,14 +2231,13 @@ def _setitem_to_parfor(self, equiv_set, loc, target, index, value, shape=None): if isinstance(value_typ, types.npytypes.Array): value_var = ir.Var(scope, mk_unique_var("$value_var"), loc) self.typemap[value_var.name] = value_typ.dtype - # Disable boundschecking, as this could be part of a slice - getitem_call = ir.Expr.getitem(value, index_var, loc, boundscheck=False) + getitem_call = ir.Expr.getitem(value, index_var, loc) self.calltypes[getitem_call] = signature( value_typ.dtype, value_typ, index_var_typ) true_block.body.append(ir.Assign(getitem_call, value_var, loc)) else: value_var = value - setitem_node = ir.SetItem(target, index_var, value_var, loc, boundscheck=False) + setitem_node = ir.SetItem(target, index_var, value_var, loc) self.calltypes[setitem_node] = signature( types.none, self.typemap[target.name], index_var_typ, el_typ) true_block.body.append(setitem_node) From 59df314c60cf064288c06e194b550495ead8468a Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Tue, 26 Nov 2019 10:09:42 -0700 Subject: [PATCH 25/33] Fix flake8 errors --- numba/cgutils.py | 22 ++++++++++++++++------ numba/config.py | 3 ++- numba/targets/arrayobj.py | 14 +++++++++----- numba/tests/test_boundscheck.py | 6 ++++++ 4 files changed, 33 insertions(+), 12 deletions(-) diff --git a/numba/cgutils.py b/numba/cgutils.py index 6c91d79d07a..ab586334568 100644 --- a/numba/cgutils.py +++ b/numba/cgutils.py @@ -657,7 +657,8 @@ def unpack_tuple(builder, tup, count=None): return vals -def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False, boundscheck=False): +def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False, + boundscheck=False): # Set boundscheck=True for any pointer access that should be # boundschecked. do_boundscheck() will handle enabling or disabling the # actual boundschecking based on the user config. @@ -667,6 +668,7 @@ def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False, bound strides=strides, layout=aryty.layout, inds=inds, wraparound=wraparound, boundscheck=boundscheck) + def do_boundscheck(context, builder, ind, dimlen, axis=None): # Boundschecking is always disabled for CUDA from .cuda.target import CUDATargetContext @@ -682,22 +684,30 @@ def _dbg(): # in the error message. if axis is not None: if isinstance(axis, int): - printf(builder, "debug: IndexError: index %d is out of bounds for axis {} with size %d\n".format(axis), ind, dimlen) + printf(builder, "debug: IndexError: index %d is out of bounds " + "for axis {} with size %d\n".format(axis), ind, dimlen) else: - printf(builder, "debug: IndexError: index %d is out of bounds for axis %d with size %d\n".format(axis), ind, axis, dimlen) + printf(builder, "debug: IndexError: index %d is out of bounds " + "for axis %d with size %d\n".format(axis), ind, axis, + dimlen) else: - printf(builder, "debug: IndexError: index %d is out of bounds for size %d\n", ind, dimlen) + printf(builder, + "debug: IndexError: index %d is out of bounds for size %d\n", + ind, dimlen) msg = "index is out of bounds" out_of_bounds_upper = builder.icmp_signed('>=', ind, dimlen) with if_unlikely(builder, out_of_bounds_upper): - if config.FULL_TRACEBACKS: _dbg() + if config.FULL_TRACEBACKS: + _dbg() context.call_conv.return_user_exc(builder, IndexError, (msg,)) out_of_bounds_lower = builder.icmp_signed('<', ind, ind.type(0)) with if_unlikely(builder, out_of_bounds_lower): - if config.FULL_TRACEBACKS: _dbg() + if config.FULL_TRACEBACKS: + _dbg() context.call_conv.return_user_exc(builder, IndexError, (msg,)) + def get_item_pointer2(context, builder, data, shape, strides, layout, inds, wraparound=False, boundscheck=False): # Set boundscheck=True for any pointer access that should be diff --git a/numba/config.py b/numba/config.py index 3764557b193..d21d696bf00 100644 --- a/numba/config.py +++ b/numba/config.py @@ -145,7 +145,8 @@ def optional_str(x): COLOR_SCHEME = _readenv("NUMBA_COLOR_SCHEME", str, "no_color") # Whether to globally enable bounds checking. The default None means - # to use the value of the flag to @njit. 0 or 1 overrides the flag globally. + # to use the value of the flag to @njit. 0 or 1 overrides the flag + # globally. BOUNDSCHECK = _readenv("NUMBA_BOUNDSCHECK", int, None) # Debug flag to control compiler debug print diff --git a/numba/targets/arrayobj.py b/numba/targets/arrayobj.py index 352ceefac47..1461116c6ba 100644 --- a/numba/targets/arrayobj.py +++ b/numba/targets/arrayobj.py @@ -315,7 +315,8 @@ def iternext_array(context, builder, sig, args, result): #------------------------------------------------------------------------------- # Basic indexing (with integers and slices only) -def basic_indexing(context, builder, aryty, ary, index_types, indices, boundscheck=True): +def basic_indexing(context, builder, aryty, ary, index_types, indices, + boundscheck=True): """ Perform basic indexing on the given array. A (data pointer, shapes, strides) tuple is returned describing @@ -1028,7 +1029,8 @@ def fancy_getitem(context, builder, sig, args, # No need to check for wraparound, as the indexers all ensure # a positive index is returned. ptr = cgutils.get_item_pointer2(context, builder, data, shapes, strides, - aryty.layout, indices, wraparound=False, boundscheck=True) + aryty.layout, indices, wraparound=False, + boundscheck=True) val = load_item(context, builder, aryty, ptr) # Since the destination is C-contiguous, no need for multi-dimensional @@ -1178,7 +1180,8 @@ def src_getitem(source_indices): with builder.if_else(use_copy, likely=False) as (if_copy, otherwise): with if_copy: builder.store( - cgutils.get_item_pointer2(context, builder, builder.load(copy_data), + cgutils.get_item_pointer2(context, builder, + builder.load(copy_data), copy_shapes, copy_strides, copy_layout, source_indices, wraparound=False), @@ -3118,8 +3121,9 @@ def _ptr_for_index(self, context, builder, arrty, arr, index): index = builder.udiv(index, shapes[dim]) indices.reverse() - ptr = cgutils.get_item_pointer2(context, builder, arr.data, shapes, - strides, arrty.layout, indices) + ptr = cgutils.get_item_pointer2(context, builder, arr.data, + shapes, strides, arrty.layout, + indices) return ptr def getitem(self, context, builder, arrty, arr, index): diff --git a/numba/tests/test_boundscheck.py b/numba/tests/test_boundscheck.py index 0a96df41981..c2f0175fbe7 100644 --- a/numba/tests/test_boundscheck.py +++ b/numba/tests/test_boundscheck.py @@ -11,17 +11,21 @@ BOUNDSCHECK_FLAGS = DEFAULT_FLAGS.copy() BOUNDSCHECK_FLAGS.set('boundscheck', True) + def basic_array_access(a): return a[10] + def slice_array_access(a): # The first index (slice) is not bounds checked return a[10:, 10] + def fancy_array_access(x): a = np.array([1, 2, 3]) return x[a] + class TestBoundsCheckNoError(MemoryLeakMixin, unittest.TestCase): def setUp(self): self.old_boundscheck = config.BOUNDSCHECK @@ -109,6 +113,7 @@ def func(): def tearDown(self): config.BOUNDSCHECK = self.old_boundscheck + # This is a separate test because the jitted functions that raise exceptions # have memory leaks. class TestBoundsCheckError(unittest.TestCase): @@ -174,5 +179,6 @@ def test_fancy_indexing_boundscheck(self): def tearDown(self): config.BOUNDSCHECK = self.old_boundscheck + if __name__ == '__main__': unittest.main() From 126efea4c3887d9195ac4fcccf28bd3be8b65fa9 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Tue, 26 Nov 2019 10:11:28 -0700 Subject: [PATCH 26/33] Update test_jit_module to know about the boundscheck flag --- numba/tests/test_jit_module.py | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/numba/tests/test_jit_module.py b/numba/tests/test_jit_module.py index 57f059adaa1..a7afa431b76 100644 --- a/numba/tests/test_jit_module.py +++ b/numba/tests/test_jit_module.py @@ -129,6 +129,7 @@ def test_jit_module_jit_options(self): jit_options = {"nopython": True, "nogil": False, "error_model": "numpy", + "boundscheck": False, } with self.create_temp_jitted_module(**jit_options) as test_module: self.assertEqual(test_module.inc.targetoptions, jit_options) @@ -148,13 +149,14 @@ def add(x, y): """ jit_options = {"nopython": True, "error_model": "numpy", + "boundscheck": False, } with self.create_temp_jitted_module(source_lines=source_lines, **jit_options) as test_module: self.assertEqual(test_module.add.targetoptions, jit_options) # Test that manual jit-wrapping overrides jit_module options self.assertEqual(test_module.inc.targetoptions, - {'nogil': True, 'forceobj': True}) + {'nogil': True, 'forceobj': True, 'boundscheck': False}) def test_jit_module_logging_output(self): logger = logging.getLogger('numba.decorators') From 4669658fdf98d916d5ef2473d8125be561fb26e7 Mon Sep 17 00:00:00 2001 From: Siu Kwan Lam Date: Wed, 4 Dec 2019 12:02:31 -0700 Subject: [PATCH 27/33] Improve the way the boundscheck config is read --- numba/cgutils.py | 9 --------- numba/cuda/compiler.py | 1 - numba/targets/arrayobj.py | 10 ++++++---- 3 files changed, 6 insertions(+), 14 deletions(-) diff --git a/numba/cgutils.py b/numba/cgutils.py index ab586334568..cdce2513163 100644 --- a/numba/cgutils.py +++ b/numba/cgutils.py @@ -670,15 +670,6 @@ def get_item_pointer(context, builder, aryty, ary, inds, wraparound=False, def do_boundscheck(context, builder, ind, dimlen, axis=None): - # Boundschecking is always disabled for CUDA - from .cuda.target import CUDATargetContext - if isinstance(context, CUDATargetContext): - return - - if (config.BOUNDSCHECK is None and not context.enable_boundscheck) \ - or config.BOUNDSCHECK == 0: - return - def _dbg(): # Remove this when we figure out how to include this information # in the error message. diff --git a/numba/cuda/compiler.py b/numba/cuda/compiler.py index a457f9451f9..fa65b2e1aa2 100644 --- a/numba/cuda/compiler.py +++ b/numba/cuda/compiler.py @@ -37,7 +37,6 @@ def compile_cuda(pyfunc, return_type, args, debug, inline): flags.set('no_compile') flags.set('no_cpython_wrapper') if debug: - flags.set('boundscheck') flags.set('debuginfo') if inline: flags.set('forceinline') diff --git a/numba/targets/arrayobj.py b/numba/targets/arrayobj.py index 1461116c6ba..b43b78eb8be 100644 --- a/numba/targets/arrayobj.py +++ b/numba/targets/arrayobj.py @@ -316,7 +316,7 @@ def iternext_array(context, builder, sig, args, result): # Basic indexing (with integers and slices only) def basic_indexing(context, builder, aryty, ary, index_types, indices, - boundscheck=True): + boundscheck=None): """ Perform basic indexing on the given array. A (data pointer, shapes, strides) tuple is returned describing @@ -400,7 +400,8 @@ def _getitem_array_generic(context, builder, return_type, aryty, ary, returning either a scalar or a view. """ dataptr, view_shapes, view_strides = \ - basic_indexing(context, builder, aryty, ary, index_types, indices) + basic_indexing(context, builder, aryty, ary, index_types, indices, + boundscheck=context.enable_boundscheck) if isinstance(return_type, types.Buffer): # Build array view @@ -478,7 +479,8 @@ def setitem_array(context, builder, sig, args): index_types, indices) try: dataptr, shapes, strides = \ - basic_indexing(context, builder, aryty, ary, index_types, indices) + basic_indexing(context, builder, aryty, ary, index_types, indices, + boundscheck=context.enable_boundscheck) except NotImplementedError: use_fancy_indexing = True else: @@ -1030,7 +1032,7 @@ def fancy_getitem(context, builder, sig, args, # a positive index is returned. ptr = cgutils.get_item_pointer2(context, builder, data, shapes, strides, aryty.layout, indices, wraparound=False, - boundscheck=True) + boundscheck=context.enable_boundscheck) val = load_item(context, builder, aryty, ptr) # Since the destination is C-contiguous, no need for multi-dimensional From 127c2d6fac30daa0a5433331ce5d65392a3ff1b5 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Wed, 4 Dec 2019 12:25:45 -0700 Subject: [PATCH 28/33] Fix the NUMBA_BOUNDSCHECK environment variable and add tests for it --- numba/targets/base.py | 14 +++++++-- numba/tests/test_boundscheck.py | 56 +++++++++++++++++++++++++++++++-- 2 files changed, 66 insertions(+), 4 deletions(-) diff --git a/numba/targets/base.py b/numba/targets/base.py index 02cc1753b9e..a7b52d69e52 100644 --- a/numba/targets/base.py +++ b/numba/targets/base.py @@ -14,7 +14,7 @@ from llvmlite.llvmpy.core import Type, Constant, LLVMException import llvmlite.binding as ll -from numba import types, utils, cgutils, typing, funcdesc, debuginfo +from numba import types, utils, cgutils, typing, funcdesc, debuginfo, config from numba import _dynfunc, _helperlib from numba.compiler_lock import global_compiler_lock from numba.pythonapi import PythonAPI @@ -197,7 +197,15 @@ class BaseContext(object): DIBuilder = debuginfo.DIBuilder # Bound checking - enable_boundscheck = False + @property + def enable_boundscheck(self): + if config.BOUNDSCHECK is not None: + return config.BOUNDSCHECK + return self._boundscheck + + @enable_boundscheck.setter + def enable_boundscheck(self, value): + self._boundscheck = value # NRT enable_nrt = False @@ -244,6 +252,8 @@ def __init__(self, typing_context): self._pid = None self._codelib_stack = [] + self._boundscheck = False + self.data_model_manager = datamodel.default_manager # Initialize diff --git a/numba/tests/test_boundscheck.py b/numba/tests/test_boundscheck.py index c2f0175fbe7..395f3f096c9 100644 --- a/numba/tests/test_boundscheck.py +++ b/numba/tests/test_boundscheck.py @@ -3,10 +3,10 @@ import numpy as np from numba.compiler import compile_isolated, DEFAULT_FLAGS -from numba import typeof, config, cuda +from numba import typeof, config, cuda, njit from numba.types import float64, intp from numba import unittest_support as unittest -from .support import MemoryLeakMixin +from .support import MemoryLeakMixin, override_env_config BOUNDSCHECK_FLAGS = DEFAULT_FLAGS.copy() BOUNDSCHECK_FLAGS.set('boundscheck', True) @@ -179,6 +179,58 @@ def test_fancy_indexing_boundscheck(self): def tearDown(self): config.BOUNDSCHECK = self.old_boundscheck +class TestBoundsEnvironmentVariable(unittest.TestCase): + def setUp(self): + self.old_boundscheck = config.BOUNDSCHECK + config.BOUNDSCHECK = None + + @njit + def default(x): + return x[1] + + @njit(boundscheck=False) + def off(x): + return x[1] + + @njit(boundscheck=True) + def on(x): + return x[1] + + self.default = default + self.off = off + self.on = on + + def test_boundscheck_unset(self): + with override_env_config('NUMBA_BOUNDSCHECK', ''): + a = np.array([1]) + + # Doesn't raise + self.default(a) + self.off(a) + + with self.assertRaises(IndexError): + self.on(a) + + def test_boundscheck_enabled(self): + with override_env_config('NUMBA_BOUNDSCHECK', '1'): + a = np.array([1]) + + with self.assertRaises(IndexError): + self.default(a) + self.off(a) + self.on(a) + + def test_boundscheck_disabled(self): + with override_env_config('NUMBA_BOUNDSCHECK', '0'): + a = np.array([1]) + + # Doesn't raise + self.default(a) + self.off(a) + self.on(a) + + def tearDown(self): + config.BOUNDSCHECK = self.old_boundscheck if __name__ == '__main__': unittest.main() From c653efb68abde2279294ee95e354be3340299823 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Wed, 4 Dec 2019 12:34:16 -0700 Subject: [PATCH 29/33] Make sure boundschecking is unconditionally disabled in cuda --- numba/cuda/target.py | 5 +++++ numba/tests/test_boundscheck.py | 10 ++++++++++ 2 files changed, 15 insertions(+) diff --git a/numba/cuda/target.py b/numba/cuda/target.py index f475551aff9..9c70ae9e1ca 100644 --- a/numba/cuda/target.py +++ b/numba/cuda/target.py @@ -60,6 +60,11 @@ class CUDATargetContext(BaseContext): strict_alignment = True DIBuilder = debuginfo.NvvmDIBuilder + @property + def enable_boundscheck(self): + # Unconditionally disabled + return False + # Overrides def create_module(self, name): return self._internal_codegen._create_empty_module(name) diff --git a/numba/tests/test_boundscheck.py b/numba/tests/test_boundscheck.py index 395f3f096c9..6f26881ed7f 100644 --- a/numba/tests/test_boundscheck.py +++ b/numba/tests/test_boundscheck.py @@ -110,6 +110,16 @@ def test_no_cuda_boundscheck(self): def func(): pass + with override_env_config('NUMBA_BOUNDSCHECK', '1'): + @cuda.jit + def func2(x, a): + a[1] = x[1] + + a = np.ones((1,)) + x = np.zeros((1,)) + # Out of bounds but doesn't raise + func2(x, a) + def tearDown(self): config.BOUNDSCHECK = self.old_boundscheck From 9a42479316f89f592bc902ed8269dbf0a6e17a1a Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Wed, 4 Dec 2019 13:40:00 -0700 Subject: [PATCH 30/33] Fix flake8 errors --- numba/tests/test_boundscheck.py | 2 ++ numba/tests/test_jit_module.py | 3 ++- 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/numba/tests/test_boundscheck.py b/numba/tests/test_boundscheck.py index 6f26881ed7f..7f57430e460 100644 --- a/numba/tests/test_boundscheck.py +++ b/numba/tests/test_boundscheck.py @@ -189,6 +189,7 @@ def test_fancy_indexing_boundscheck(self): def tearDown(self): config.BOUNDSCHECK = self.old_boundscheck + class TestBoundsEnvironmentVariable(unittest.TestCase): def setUp(self): self.old_boundscheck = config.BOUNDSCHECK @@ -242,5 +243,6 @@ def test_boundscheck_disabled(self): def tearDown(self): config.BOUNDSCHECK = self.old_boundscheck + if __name__ == '__main__': unittest.main() diff --git a/numba/tests/test_jit_module.py b/numba/tests/test_jit_module.py index a7afa431b76..05e74d2ae9e 100644 --- a/numba/tests/test_jit_module.py +++ b/numba/tests/test_jit_module.py @@ -156,7 +156,8 @@ def add(x, y): self.assertEqual(test_module.add.targetoptions, jit_options) # Test that manual jit-wrapping overrides jit_module options self.assertEqual(test_module.inc.targetoptions, - {'nogil': True, 'forceobj': True, 'boundscheck': False}) + {'nogil': True, 'forceobj': True, + 'boundscheck': False}) def test_jit_module_logging_output(self): logger = logging.getLogger('numba.decorators') From e55e804ba00ca1eef7d30bbb287c9a66a1233951 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Wed, 4 Dec 2019 14:23:17 -0700 Subject: [PATCH 31/33] Fix the boundscheck tests when NUMBA_ENABLE_CUDASIM=1 --- numba/cuda/simulator/api.py | 8 +++++++- numba/tests/test_boundscheck.py | 6 ++++-- 2 files changed, 11 insertions(+), 3 deletions(-) diff --git a/numba/cuda/simulator/api.py b/numba/cuda/simulator/api.py index e3adfbdc224..2c2d82334b4 100644 --- a/numba/cuda/simulator/api.py +++ b/numba/cuda/simulator/api.py @@ -73,7 +73,13 @@ def elapsed_time(self, event): def jit(func_or_sig=None, device=False, debug=False, argtypes=None, - inline=False, restype=None, fastmath=False, link=None): + inline=False, restype=None, fastmath=False, link=None, + boundscheck=None, + ): + # Here for API compatibility + if boundscheck is not None: + raise NotImplementedError("bounds checking is not supported for CUDA") + if link is not None: raise NotImplementedError('Cannot link PTX in the simulator') # Check for first argument specifying types - in that case the diff --git a/numba/tests/test_boundscheck.py b/numba/tests/test_boundscheck.py index 7f57430e460..bd6895c65fe 100644 --- a/numba/tests/test_boundscheck.py +++ b/numba/tests/test_boundscheck.py @@ -117,8 +117,10 @@ def func2(x, a): a = np.ones((1,)) x = np.zeros((1,)) - # Out of bounds but doesn't raise - func2(x, a) + # Out of bounds but doesn't raise (it does raise in the simulator, + # so skip there) + if not config.ENABLE_CUDASIM: + func2(x, a) def tearDown(self): config.BOUNDSCHECK = self.old_boundscheck From a0140c22b5cdf69182b151be8f741fbc0e5be395 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Wed, 4 Dec 2019 14:57:40 -0700 Subject: [PATCH 32/33] Skip the CUDA boundscheck test if CUDA isn't available --- numba/tests/test_boundscheck.py | 1 + 1 file changed, 1 insertion(+) diff --git a/numba/tests/test_boundscheck.py b/numba/tests/test_boundscheck.py index bd6895c65fe..a3ebf7a9c2a 100644 --- a/numba/tests/test_boundscheck.py +++ b/numba/tests/test_boundscheck.py @@ -104,6 +104,7 @@ def test_fancy_indexing_boundscheck(self): # Doesn't raise boundscheck(b) + @unittest.skipIf(not cuda.is_available(), "NO CUDA") def test_no_cuda_boundscheck(self): with self.assertRaises(NotImplementedError): @cuda.jit(boundscheck=True) From 129459d857eb25f2136ffde4a6d1f38cb80bc492 Mon Sep 17 00:00:00 2001 From: Aaron Meurer Date: Wed, 4 Dec 2019 15:40:37 -0700 Subject: [PATCH 33/33] Use the dtype directly in the test Hopefully this fill fix the test failures on Windows. --- numba/tests/test_boundscheck.py | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/numba/tests/test_boundscheck.py b/numba/tests/test_boundscheck.py index a3ebf7a9c2a..8ee117a8ffa 100644 --- a/numba/tests/test_boundscheck.py +++ b/numba/tests/test_boundscheck.py @@ -4,7 +4,7 @@ from numba.compiler import compile_isolated, DEFAULT_FLAGS from numba import typeof, config, cuda, njit -from numba.types import float64, intp +from numba.types import float64 from numba import unittest_support as unittest from .support import MemoryLeakMixin, override_env_config @@ -87,7 +87,7 @@ def test_fancy_indexing_boundscheck(self): fancy_array_access(b) at = typeof(a) - rt = intp[:] + rt = at.dtype[:] c_noboundscheck = compile_isolated(fancy_array_access, [at], return_type=rt, flags=DEFAULT_FLAGS) @@ -181,7 +181,7 @@ def test_fancy_indexing_boundscheck(self): fancy_array_access(b) at = typeof(a) - rt = intp[:] + rt = at.dtype[:] c_boundscheck = compile_isolated(fancy_array_access, [at], return_type=rt, flags=BOUNDSCHECK_FLAGS)