Fix few small issues with filters on large images
@@ -412,13 +412,15 @@ def _check_nd_args(input, weights, mode, origins, wghts_name='filter weights'):
raise TypeError('Complex type not supported.')
_check_mode(mode)
# The integer type to use for indices in the input array
- # We will always assume that weights will always use int32 however
# The indices actually use byte positions and we can't just use
# input.nbytes since that won't tell us the number of bytes between the
# first and last elements when the array is non-contiguous
- nbytes = sum((x-1)*strd for x, strd in zip(input.shape, input.strides))
- nbytes += input.dtype.itemsize
+ nbytes = sum((x-1)*abs(stride) for x, stride in
+ zip(input.shape, input.strides)) + input.dtype.itemsize
int_type = 'int' if nbytes < (1 << 31) else 'ptrdiff_t'
+ # However, weights must always be 2 GiB or less
+ if weights.nbytes > (1 << 31):
+ raise RuntimeError('weights must be 2 GiB or less, use FFTs instead')
weight_dims = [x for x in weights.shape if x != 0]
if len(weight_dims) != input.ndim:
raise RuntimeError('{} array has incorrect shape.'.format(wghts_name))
@@ -564,7 +566,7 @@ def _generate_nd_kernel(name, pre, found, post, mode, wshape, int_type,
name = 'cupy_ndimage_{}_{}d_{}_w{}'.format(
name, ndim, mode, '_'.join(['{}'.format(j) for j in wshape]))
- if int_type == 'size_t':
+ if int_type == 'ptrdiff_t':
name += '_i64'
return cupy.ElementwiseKernel(in_params, out_params, operation, name,
reduce_dims=False, preamble=preamble,
