def dotc(x, y, out=None): """Computes the dot product of x.conj() and y.""" dtype = x.dtype.char if dtype in 'fd': return dot(x, y, out=out) elif dtype == 'F': func = cublas.cdotc elif dtype == 'D': func = cublas.zdotc else: raise TypeError('invalid dtype') _check_two_vectors(x, y) handle = device.get_cublas_handle() result_dtype = dtype result_ptr, result, orig_mode = _setup_result_ptr( handle, out, result_dtype) try: func(handle, x.size, x.data.ptr, 1, y.data.ptr, 1, result_ptr) finally: cublas.setPointerMode(handle, orig_mode) if out is None: out = result elif out.dtype != result_dtype: _core.elementwise_copy(result, out) return out
def _iamaxmin(x, out, name): if x.ndim != 1: raise ValueError('x must be a 1D array (actual: {})'.format(x.ndim)) dtype = x.dtype.char if dtype == 'f': t = 's' elif dtype == 'd': t = 'd' elif dtype == 'F': t = 'c' elif dtype == 'D': t = 'z' else: raise TypeError('invalid dtype') func = getattr(cublas, 'i' + t + name) handle = device.get_cublas_handle() result_dtype = 'i' result_ptr, result, orig_mode = _setup_result_ptr( handle, out, result_dtype) try: func(handle, x.size, x.data.ptr, 1, result_ptr) finally: cublas.setPointerMode(handle, orig_mode) if out is None: out = result elif out.dtype != result_dtype: _core.elementwise_copy(result, out) return out
def _run_1d_filters(filters, input, args, output, mode, cval, origin=0): """ Runs a series of 1D filters forming an nd filter. The filters must be a list of callables that take input, arg, axis, output, mode, cval, origin. The args is a list of values that are passed for the arg value to the filter. Individual filters can be None causing that axis to be skipped. """ output_orig = output output = _util._get_output(output, input) modes = _util._fix_sequence_arg(mode, input.ndim, 'mode', _util._check_mode) # for filters, "wrap" is a synonym for "grid-wrap". modes = ['grid-wrap' if m == 'wrap' else m for m in modes] origins = _util._fix_sequence_arg(origin, input.ndim, 'origin', int) n_filters = sum(filter is not None for filter in filters) if n_filters == 0: _core.elementwise_copy(input, output) return output # We can't operate in-place efficiently, so use a 2-buffer system temp = _util._get_output(output.dtype, input) if n_filters > 1 else None first = True iterator = zip(filters, args, modes, origins) for axis, (fltr, arg, mode, origin) in enumerate(iterator): if fltr is None: continue fltr(input, arg, axis, output, mode, cval, origin) input, output = output, temp if first else input first = False if isinstance(output_orig, cupy.ndarray) and input is not output_orig: _core.elementwise_copy(input, output_orig) input = output_orig return input
def nrm2(x, out=None): """Computes the Euclidean norm of vector x.""" if x.ndim != 1: raise ValueError('x must be a 1D array (actual: {})'.format(x.ndim)) dtype = x.dtype.char if dtype == 'f': func = cublas.snrm2 elif dtype == 'd': func = cublas.dnrm2 elif dtype == 'F': func = cublas.scnrm2 elif dtype == 'D': func = cublas.dznrm2 else: raise TypeError('invalid dtype') handle = device.get_cublas_handle() result_dtype = dtype.lower() result_ptr, result, orig_mode = _setup_result_ptr( handle, out, result_dtype) try: func(handle, x.size, x.data.ptr, 1, result_ptr) finally: cublas.setPointerMode(handle, orig_mode) if out is None: out = result elif out.dtype != result_dtype: _core.elementwise_copy(result, out) return out
def gerc(alpha, x, y, a): """Computes a += alpha * x @ y.T.conj() Note: ''a'' will be updated. """ dtype = a.dtype.char if dtype in 'fd': return ger(alpha, x, y, a) elif dtype == 'F': func = cublas.cgerc elif dtype == 'D': func = cublas.zgerc else: raise TypeError('invalid dtype') assert a.ndim == 2 assert x.ndim == y.ndim == 1 assert a.dtype == x.dtype == y.dtype m, n = a.shape assert x.shape[0] == m assert y.shape[0] == n handle = device.get_cublas_handle() alpha, alpha_ptr, orig_mode = _setup_scalar_ptr(handle, alpha, dtype) x_ptr, y_ptr = x.data.ptr, y.data.ptr try: if a._f_contiguous: func(handle, m, n, alpha_ptr, x_ptr, 1, y_ptr, 1, a.data.ptr, m) else: aa = a.copy(order='F') func(handle, m, n, alpha_ptr, x_ptr, 1, y_ptr, 1, aa.data.ptr, m) _core.elementwise_copy(aa, a) finally: cublas.setPointerMode(handle, orig_mode)
def check_copy(self, dtype, src_id, dst_id): with cuda.Device(src_id): src = testing.shaped_arange((2, 3, 4), dtype=dtype) with cuda.Device(dst_id): dst = cupy.empty((2, 3, 4), dtype=dtype) _core.elementwise_copy(src, dst) testing.assert_allclose(src, dst)
def copyto(dst, src, casting='same_kind', where=None): """Copies values from one array to another with broadcasting. This function can be called for arrays on different devices. In this case, casting, ``where``, and broadcasting is not supported, and an exception is raised if these are used. Args: dst (cupy.ndarray): Target array. src (cupy.ndarray): Source array. casting (str): Casting rule. See :func:`numpy.can_cast` for detail. where (cupy.ndarray of bool): If specified, this array acts as a mask, and an element is copied only if the corresponding element of ``where`` is True. .. seealso:: :func:`numpy.copyto` """ src_type = type(src) src_is_python_scalar = src_type in (int, bool, float, complex, fusion._FusionVarScalar, _fusion_interface._ScalarProxy) if src_is_python_scalar: src_dtype = numpy.dtype(type(src)) can_cast = numpy.can_cast(src, dst.dtype, casting) else: src_dtype = src.dtype can_cast = numpy.can_cast(src_dtype, dst.dtype, casting) if not can_cast: raise TypeError('Cannot cast %s to %s in %s casting mode' % (src_dtype, dst.dtype, casting)) if fusion._is_fusing(): if where is None: _core.elementwise_copy(src, dst) else: fusion._call_ufunc(search._where_ufunc, where, src, dst, dst) return if dst.size == 0: return if src_is_python_scalar and where is None: dst.fill(src) return if where is None: if _can_memcpy(dst, src): dst.data.copy_from_async(src.data, src.nbytes) else: device = dst.device with device: if src.device != device: src = src.copy() _core.elementwise_copy(src, dst) else: _core.elementwise_copy_where(src, where, dst)
def dgmm(side, a, x, out=None, incx=1): """Computes diag(x) @ a or a @ diag(x) Computes diag(x) @ a if side is 'L', a @ diag(x) if side is 'R'. """ assert a.ndim == 2 assert 0 <= x.ndim <= 2 assert a.dtype == x.dtype dtype = a.dtype.char if dtype == 'f': func = cublas.sdgmm elif dtype == 'd': func = cublas.ddgmm elif dtype == 'F': func = cublas.cdgmm elif dtype == 'D': func = cublas.zdgmm else: raise TypeError('invalid dtype') if side == 'L' or side == cublas.CUBLAS_SIDE_LEFT: side = cublas.CUBLAS_SIDE_LEFT elif side == 'R' or side == cublas.CUBLAS_SIDE_RIGHT: side = cublas.CUBLAS_SIDE_RIGHT else: raise ValueError('invalid side (actual: {})'.format(side)) m, n = a.shape if side == cublas.CUBLAS_SIDE_LEFT: assert x.size >= (m - 1) * abs(incx) + 1 else: assert x.size >= (n - 1) * abs(incx) + 1 if out is None: if a._c_contiguous: order = 'C' else: order = 'F' out = cupy.empty((m, n), dtype=dtype, order=order) else: assert out.ndim == 2 assert out.shape == a.shape assert out.dtype == a.dtype handle = device.get_cublas_handle() if out._c_contiguous: if not a._c_contiguous: a = a.copy(order='C') func(handle, 1 - side, n, m, a.data.ptr, n, x.data.ptr, incx, out.data.ptr, n) else: if not a._f_contiguous: a = a.copy(order='F') c = out if not out._f_contiguous: c = out.copy(order='F') func(handle, side, m, n, a.data.ptr, m, x.data.ptr, incx, c.data.ptr, m) if not out._f_contiguous: _core.elementwise_copy(c, out) return out
def fourier_gaussian(input, sigma, n=-1, axis=-1, output=None): """Multidimensional Gaussian shift filter. The array is multiplied with the Fourier transform of a (separable) Gaussian kernel. Args: input (cupy.ndarray): The input array. sigma (float or sequence of float): The sigma of the Gaussian kernel. If a float, `sigma` is the same for all axes. If a sequence, `sigma` has to contain one value for each axis. n (int, optional): If `n` is negative (default), then the input is assumed to be the result of a complex fft. If `n` is larger than or equal to zero, the input is assumed to be the result of a real fft, and `n` gives the length of the array before transformation along the real transform direction. axis (int, optional): The axis of the real transform (only used when ``n > -1``). output (cupy.ndarray, optional): If given, the result of shifting the input is placed in this array. Returns: output (cupy.ndarray): The filtered output. """ ndim = input.ndim output = _get_output_fourier(output, input) axis = internal._normalize_axis_index(axis, ndim) sigmas = _util._fix_sequence_arg(sigma, ndim, 'sigma') _core.elementwise_copy(input, output) for ax, (sigmak, ax_size) in enumerate(zip(sigmas, output.shape)): # compute the frequency grid in Hz if ax == axis and n > 0: arr = cupy.arange(ax_size, dtype=output.real.dtype) arr /= n else: arr = cupy.fft.fftfreq(ax_size) arr = arr.astype(output.real.dtype, copy=False) # compute the Gaussian weights arr *= arr scale = sigmak * sigmak / -2 arr *= (4 * numpy.pi * numpy.pi) * scale cupy.exp(arr, out=arr) # reshape for broadcasting arr = _reshape_nd(arr, ndim=ndim, axis=ax) output *= arr return output
def tile(A, reps): """Construct an array by repeating A the number of times given by reps. Args: A (cupy.ndarray): Array to transform. reps (int or tuple): The number of repeats. Returns: cupy.ndarray: Transformed array with repeats. .. seealso:: :func:`numpy.tile` """ try: tup = tuple(reps) except TypeError: tup = (reps,) d = len(tup) if tup.count(1) == len(tup) and isinstance(A, cupy.ndarray): # Fixes the problem that the function does not make a copy if A is a # array and the repetitions are 1 in all dimensions return cupy.array(A, copy=True, ndmin=d) else: # Note that no copy of zero-sized arrays is made. However since they # have no data there is no risk of an inadvertent overwrite. c = cupy.array(A, copy=False, ndmin=d) if d < c.ndim: tup = (1,) * (c.ndim - d) + tup shape_out = tuple(s * t for s, t in zip(c.shape, tup)) if c.size == 0: return cupy.empty(shape_out, dtype=c.dtype) c_shape = [] ret_shape = [] for dim_in, nrep in zip(c.shape, tup): if nrep == 1: c_shape.append(dim_in) ret_shape.append(dim_in) elif dim_in == 1: c_shape.append(dim_in) ret_shape.append(nrep) else: c_shape.append(1) c_shape.append(dim_in) ret_shape.append(nrep) ret_shape.append(dim_in) ret = cupy.empty(ret_shape, dtype=c.dtype) if ret.size: _core.elementwise_copy(c.reshape(c_shape), ret) return ret.reshape(shape_out)
def fourier_shift(input, shift, n=-1, axis=-1, output=None): """Multidimensional Fourier shift filter. The array is multiplied with the Fourier transform of a shift operation. Args: input (cupy.ndarray): The input array. This should be in the Fourier domain. shift (float or sequence of float): The size of shift. If a float, `shift` is the same for all axes. If a sequence, `shift` has to contain one value for each axis. n (int, optional): If `n` is negative (default), then the input is assumed to be the result of a complex fft. If `n` is larger than or equal to zero, the input is assumed to be the result of a real fft, and `n` gives the length of the array before transformation along the real transform direction. axis (int, optional): The axis of the real transform (only used when ``n > -1``). output (cupy.ndarray, optional): If given, the result of shifting the input is placed in this array. Returns: output (cupy.ndarray): The shifted output (in the Fourier domain). """ ndim = input.ndim output = _get_output_fourier(output, input, complex_only=True) axis = internal._normalize_axis_index(axis, ndim) shifts = _util._fix_sequence_arg(shift, ndim, 'shift') _core.elementwise_copy(input, output) for ax, (shiftk, ax_size) in enumerate(zip(shifts, output.shape)): if shiftk == 0: continue if ax == axis and n > 0: # cp.fft.rfftfreq(ax_size) * (-2j * numpy.pi * shiftk * ax_size/n) arr = cupy.arange(ax_size, dtype=output.dtype) arr *= -2j * numpy.pi * shiftk / n else: arr = cupy.fft.fftfreq(ax_size) arr = arr * (-2j * numpy.pi * shiftk) cupy.exp(arr, out=arr) # reshape for broadcasting arr = _reshape_nd(arr, ndim=ndim, axis=ax) output *= arr return output
def sum(self, axis=None, dtype=None, out=None): """Sums the matrix elements over a given axis. Args: axis (int or ``None``): Axis along which the sum is comuted. If it is ``None``, it computes the sum of all the elements. Select from ``{None, 0, 1, -2, -1}``. dtype: The type of returned matrix. If it is not specified, type of the array is used. out (cupy.ndarray): Output matrix. Returns: cupy.ndarray: Summed array. .. seealso:: :meth:`scipy.sparse.spmatrix.sum` """ _sputils.validateaxis(axis) # This implementation uses multiplication, though it is not efficient # for some matrix types. These should override this function. m, n = self.shape if axis is None: return self.dot(cupy.ones(n, dtype=self.dtype)).sum(dtype=dtype, out=out) if axis < 0: axis += 2 if axis == 0: ret = self.T.dot(cupy.ones(m, dtype=self.dtype)).reshape(1, n) else: # axis == 1 ret = self.dot(cupy.ones(n, dtype=self.dtype)).reshape(m, 1) if out is not None: if out.shape != ret.shape: raise ValueError('dimensions do not match') _core.elementwise_copy(ret, out) return out elif dtype is not None: return ret.astype(dtype, copy=False) else: return ret
def _call_kernel(kernel, input, weights, output, structure=None, weights_dtype=numpy.float64, structure_dtype=numpy.float64): """ Calls a constructed ElementwiseKernel. The kernel must take an input image, an optional array of weights, an optional array for the structure, and an output array. weights and structure can be given as None (structure defaults to None) in which case they are not passed to the kernel at all. If the output is given as None then it will be allocated in this function. This function deals with making sure that the weights and structure are contiguous and float64 (or bool for weights that are footprints)*, that the output is allocated and appriopately shaped. This also deals with the situation that the input and output arrays overlap in memory. * weights is always cast to float64 or bool in order to get an output compatible with SciPy, though float32 might be sufficient when input dtype is low precision. If weights_dtype is passed as weights.dtype then no dtype conversion will occur. The input and output are never converted. """ args = [input] complex_output = input.dtype.kind == 'c' if weights is not None: weights = cupy.ascontiguousarray(weights, weights_dtype) complex_output = complex_output or weights.dtype.kind == 'c' args.append(weights) if structure is not None: structure = cupy.ascontiguousarray(structure, structure_dtype) args.append(structure) output = _util._get_output(output, input, None, complex_output) needs_temp = cupy.shares_memory(output, input, 'MAY_SHARE_BOUNDS') if needs_temp: output, temp = _util._get_output(output.dtype, input), output args.append(output) kernel(*args) if needs_temp: _core.elementwise_copy(temp, output) output = temp return output
def spline_filter(input, order=3, output=cupy.float64, mode='mirror'): """Multidimensional spline filter. Args: input (cupy.ndarray): The input array. order (int): The order of the spline interpolation, default is 3. Must be in the range 0-5. output (cupy.ndarray or dtype, optional): The array in which to place the output, or the dtype of the returned array. Default is ``numpy.float64``. mode (str): Points outside the boundaries of the input are filled according to the given mode (``'constant'``, ``'nearest'``, ``'mirror'``, ``'reflect'``, ``'wrap'``, ``'grid-mirror'``, ``'grid-wrap'``, ``'grid-constant'`` or ``'opencv'``). Returns: cupy.ndarray: The result of prefiltering the input. .. seealso:: :func:`scipy.spline_filter1d` """ if order < 2 or order > 5: raise RuntimeError('spline order not supported') x = input temp, data_dtype, output_dtype = _get_spline_output(x, output) if order not in [0, 1] and input.ndim > 0: for axis in range(x.ndim): spline_filter1d(x, order, axis, output=temp, mode=mode) x = temp if isinstance(output, cupy.ndarray): _core.elementwise_copy(temp, output) else: output = temp if output.dtype != output_dtype: output = output.astype(output_dtype) return output
def _get_spline_output(input, output): """Create workspace array, temp, and the final dtype for the output. Differs from SciPy by not always forcing the internal floating point dtype to be double precision. """ complex_data = input.dtype.kind == 'c' if complex_data: min_float_dtype = cupy.complex64 else: min_float_dtype = cupy.float32 if isinstance(output, cupy.ndarray): if complex_data and output.dtype.kind != 'c': raise ValueError( 'output must have complex dtype for complex inputs') float_dtype = cupy.promote_types(output.dtype, min_float_dtype) output_dtype = output.dtype else: if output is None: output = output_dtype = input.dtype else: output_dtype = cupy.dtype(output) float_dtype = cupy.promote_types(output, min_float_dtype) if (isinstance(output, cupy.ndarray) and output.dtype == float_dtype == output_dtype and output.flags.c_contiguous): if output is not input: _core.elementwise_copy(input, output) temp = output else: temp = input.astype(float_dtype, copy=False) temp = cupy.ascontiguousarray(temp) if cupy.shares_memory(temp, input, 'MAY_SHARE_BOUNDS'): temp = temp.copy() return temp, float_dtype, output_dtype
def copyto(dst, src, casting='same_kind', where=None): """Copies values from one array to another with broadcasting. This function can be called for arrays on different devices. In this case, casting, ``where``, and broadcasting is not supported, and an exception is raised if these are used. Args: dst (cupy.ndarray): Target array. src (cupy.ndarray): Source array. casting (str): Casting rule. See :func:`numpy.can_cast` for detail. where (cupy.ndarray of bool): If specified, this array acts as a mask, and an element is copied only if the corresponding element of ``where`` is True. .. seealso:: :func:`numpy.copyto` """ src_is_numpy_scalar = False src_type = type(src) src_is_python_scalar = src_type in ( int, bool, float, complex, fusion._FusionVarScalar, _fusion_interface._ScalarProxy) if src_is_python_scalar: src_dtype = numpy.dtype(type(src)) can_cast = numpy.can_cast(src, dst.dtype, casting) elif isinstance(src, numpy.ndarray) or numpy.isscalar(src): if src.size != 1: raise ValueError( 'non-scalar numpy.ndarray cannot be used for copyto') src_dtype = src.dtype can_cast = numpy.can_cast(src, dst.dtype, casting) src = src.item() src_is_numpy_scalar = True else: src_dtype = src.dtype can_cast = numpy.can_cast(src_dtype, dst.dtype, casting) if not can_cast: raise TypeError('Cannot cast %s to %s in %s casting mode' % (src_dtype, dst.dtype, casting)) if fusion._is_fusing(): # TODO(kataoka): NumPy allows stripping leading unit dimensions. # But fusion array proxy does not currently support # `shape` and `squeeze`. if where is None: _core.elementwise_copy(src, dst) else: fusion._call_ufunc(search._where_ufunc, where, src, dst, dst) return if not src_is_python_scalar and not src_is_numpy_scalar: # Check broadcast condition # - for fast-paths and # - for a better error message (than ufunc's). # NumPy allows stripping leading unit dimensions. if not all([ s in (d, 1) for s, d in itertools.zip_longest( reversed(src.shape), reversed(dst.shape), fillvalue=1) ]): raise ValueError( "could not broadcast input array " f"from shape {src.shape} into shape {dst.shape}") squeeze_ndim = src.ndim - dst.ndim if squeeze_ndim > 0: # always succeeds because broadcast conition is checked. src = src.squeeze(tuple(range(squeeze_ndim))) if where is not None: _core.elementwise_copy(src, dst, _where=where) return if dst.size == 0: return if src_is_python_scalar or src_is_numpy_scalar: _core.elementwise_copy(src, dst) return if _can_memcpy(dst, src): dst.data.copy_from_async(src.data, src.nbytes) return device = dst.device prev_device = runtime.getDevice() try: runtime.setDevice(device.id) if src.device != device: src = src.copy() _core.elementwise_copy(src, dst) finally: runtime.setDevice(prev_device)
def label(input, structure=None, output=None): """Labels features in an array. Args: input (cupy.ndarray): The input array. structure (array_like or None): A structuring element that defines feature connections. ```structure``` must be centersymmetric. If None, structure is automatically generated with a squared connectivity equal to one. output (cupy.ndarray, dtype or None): The array in which to place the output. Returns: label (cupy.ndarray): An integer array where each unique feature in ```input``` has a unique label in the array. num_features (int): Number of features found. .. warning:: This function may synchronize the device. .. seealso:: :func:`scipy.ndimage.label` """ if not isinstance(input, cupy.ndarray): raise TypeError('input must be cupy.ndarray') if input.dtype.char in 'FD': raise TypeError('Complex type not supported') if structure is None: structure = _generate_binary_structure(input.ndim, 1) elif isinstance(structure, cupy.ndarray): structure = cupy.asnumpy(structure) structure = numpy.array(structure, dtype=bool) if structure.ndim != input.ndim: raise RuntimeError('structure and input must have equal rank') for i in structure.shape: if i != 3: raise ValueError('structure dimensions must be equal to 3') if isinstance(output, cupy.ndarray): if output.shape != input.shape: raise ValueError("output shape not correct") caller_provided_output = True else: caller_provided_output = False if output is None: output = cupy.empty(input.shape, numpy.int32) else: output = cupy.empty(input.shape, output) if input.size == 0: # empty maxlabel = 0 elif input.ndim == 0: # 0-dim array maxlabel = 0 if input.item() == 0 else 1 output.fill(maxlabel) else: if output.dtype != numpy.int32: y = cupy.empty(input.shape, numpy.int32) else: y = output maxlabel = _label(input, structure, y) if output.dtype != numpy.int32: _core.elementwise_copy(y, output) if caller_provided_output: return maxlabel else: return output, maxlabel
def spline_filter1d(input, order=3, axis=-1, output=cupy.float64, mode='mirror'): """ Calculate a 1-D spline filter along the given axis. The lines of the array along the given axis are filtered by a spline filter. The order of the spline must be >= 2 and <= 5. Args: input (cupy.ndarray): The input array. order (int): The order of the spline interpolation, default is 3. Must be in the range 0-5. axis (int): The axis along which the spline filter is applied. Default is the last axis. output (cupy.ndarray or dtype, optional): The array in which to place the output, or the dtype of the returned array. Default is ``numpy.float64``. mode (str): Points outside the boundaries of the input are filled according to the given mode (``'constant'``, ``'nearest'``, ``'mirror'``, ``'reflect'``, ``'wrap'``, ``'grid-mirror'``, ``'grid-wrap'``, ``'grid-constant'`` or ``'opencv'``). Returns: cupy.ndarray: The result of prefiltering the input. .. seealso:: :func:`scipy.spline_filter1d` """ if order < 0 or order > 5: raise RuntimeError('spline order not supported') x = input ndim = x.ndim axis = internal._normalize_axis_index(axis, ndim) # order 0, 1 don't require reshaping as no CUDA kernel will be called # scalar or size 1 arrays also don't need to be filtered run_kernel = not (order < 2 or x.ndim == 0 or x.shape[axis] == 1) if not run_kernel: output = _util._get_output(output, input) _core.elementwise_copy(x, output) return output temp, data_dtype, output_dtype = _get_spline_output(x, output) data_type = cupy._core._scalar.get_typename(temp.dtype) pole_type = cupy._core._scalar.get_typename(temp.real.dtype) index_type = _util._get_inttype(input) index_dtype = cupy.int32 if index_type == 'int' else cupy.int64 n_samples = x.shape[axis] n_signals = x.size // n_samples info = cupy.array((n_signals, n_samples) + x.shape, dtype=index_dtype) # empirical choice of block size that seemed to work well block_size = max(2**math.ceil(numpy.log2(n_samples / 32)), 8) kern = _spline_prefilter_core.get_raw_spline1d_kernel( axis, ndim, mode, order=order, index_type=index_type, data_type=data_type, pole_type=pole_type, block_size=block_size, ) # Due to recursive nature, a given line of data must be processed by a # single thread. n_signals lines will be processed in total. block = (block_size, ) grid = ((n_signals + block[0] - 1) // block[0], ) # apply prefilter gain poles = _spline_prefilter_core.get_poles(order=order) temp *= _spline_prefilter_core.get_gain(poles) # apply caual + anti-causal IIR spline filters kern(grid, block, (temp, info)) if isinstance(output, cupy.ndarray) and temp is not output: # copy kernel output into the user-provided output array _core.elementwise_copy(temp, output) return output return temp.astype(output_dtype, copy=False)
def geam(transa, transb, alpha, a, beta, b, out=None): """Computes alpha * op(a) + beta * op(b) op(a) = a if transa is 'N', op(a) = a.T if transa is 'T', op(a) = a.T.conj() if transa is 'H'. op(b) = b if transb is 'N', op(b) = b.T if transb is 'T', op(b) = b.T.conj() if transb is 'H'. """ assert a.ndim == b.ndim == 2 assert a.dtype == b.dtype dtype = a.dtype.char if dtype == 'f': func = cublas.sgeam elif dtype == 'd': func = cublas.dgeam elif dtype == 'F': func = cublas.cgeam elif dtype == 'D': func = cublas.zgeam else: raise TypeError('invalid dtype') transa = _trans_to_cublas_op(transa) transb = _trans_to_cublas_op(transb) if transa == cublas.CUBLAS_OP_N: m, n = a.shape else: n, m = a.shape if transb == cublas.CUBLAS_OP_N: assert b.shape == (m, n) else: assert b.shape == (n, m) if out is None: out = cupy.empty((m, n), dtype=dtype, order='F') else: assert out.ndim == 2 assert out.shape == (m, n) assert out.dtype == dtype alpha, alpha_ptr = _get_scalar_ptr(alpha, a.dtype) beta, beta_ptr = _get_scalar_ptr(beta, a.dtype) handle = device.get_cublas_handle() orig_mode = cublas.getPointerMode(handle) if isinstance(alpha, cupy.ndarray) or isinstance(beta, cupy.ndarray): if not isinstance(alpha, cupy.ndarray): alpha = cupy.array(alpha) alpha_ptr = alpha.data.ptr if not isinstance(beta, cupy.ndarray): beta = cupy.array(beta) beta_ptr = beta.data.ptr cublas.setPointerMode(handle, cublas.CUBLAS_POINTER_MODE_DEVICE) else: cublas.setPointerMode(handle, cublas.CUBLAS_POINTER_MODE_HOST) lda, transa = _decide_ld_and_trans(a, transa) ldb, transb = _decide_ld_and_trans(b, transb) if not (lda is None or ldb is None): if out._f_contiguous: try: func(handle, transa, transb, m, n, alpha_ptr, a.data.ptr, lda, beta_ptr, b.data.ptr, ldb, out.data.ptr, m) finally: cublas.setPointerMode(handle, orig_mode) return out elif out._c_contiguous: # Computes alpha * a.T + beta * b.T try: func(handle, 1-transa, 1-transb, n, m, alpha_ptr, a.data.ptr, lda, beta_ptr, b.data.ptr, ldb, out.data.ptr, n) finally: cublas.setPointerMode(handle, orig_mode) return out a, lda = _change_order_if_necessary(a, lda) b, ldb = _change_order_if_necessary(b, ldb) c = out if not out._f_contiguous: c = out.copy(order='F') try: func(handle, transa, transb, m, n, alpha_ptr, a.data.ptr, lda, beta_ptr, b.data.ptr, ldb, c.data.ptr, m) finally: cublas.setPointerMode(handle, orig_mode) if not out._f_contiguous: _core.elementwise_copy(c, out) return out
def _binary_erosion(input, structure, iterations, mask, output, border_value, origin, invert, brute_force=True): try: iterations = operator.index(iterations) except TypeError: raise TypeError('iterations parameter should be an integer') if input.dtype.kind == 'c': raise TypeError('Complex type not supported') if structure is None: structure = generate_binary_structure(input.ndim, 1) all_weights_nonzero = input.ndim == 1 center_is_true = True default_structure = True else: structure = structure.astype(dtype=bool, copy=False) # transfer to CPU for use in determining if it is fully dense # structure_cpu = cupy.asnumpy(structure) default_structure = False if structure.ndim != input.ndim: raise RuntimeError('structure and input must have same dimensionality') if not structure.flags.c_contiguous: structure = cupy.ascontiguousarray(structure) if structure.size < 1: raise RuntimeError('structure must not be empty') if mask is not None: if mask.shape != input.shape: raise RuntimeError('mask and input must have equal sizes') if not mask.flags.c_contiguous: mask = cupy.ascontiguousarray(mask) masked = True else: masked = False origin = _util._fix_sequence_arg(origin, input.ndim, 'origin', int) if isinstance(output, cupy.ndarray): if output.dtype.kind == 'c': raise TypeError('Complex output type not supported') else: output = bool output = _util._get_output(output, input) temp_needed = cupy.shares_memory(output, input, 'MAY_SHARE_BOUNDS') if temp_needed: # input and output arrays cannot share memory temp = output output = _util._get_output(output.dtype, input) if structure.ndim == 0: # kernel doesn't handle ndim=0, so special case it here if float(structure): output[...] = cupy.asarray(input, dtype=bool) else: output[...] = ~cupy.asarray(input, dtype=bool) return output origin = tuple(origin) int_type = _util._get_inttype(input) offsets = _filters_core._origins_to_offsets(origin, structure.shape) if not default_structure: # synchronize required to determine if all weights are non-zero nnz = int(cupy.count_nonzero(structure)) all_weights_nonzero = nnz == structure.size if all_weights_nonzero: center_is_true = True else: center_is_true = _center_is_true(structure, origin) erode_kernel = _get_binary_erosion_kernel( structure.shape, int_type, offsets, center_is_true, border_value, invert, masked, all_weights_nonzero, ) if iterations == 1: if masked: output = erode_kernel(input, structure, mask, output) else: output = erode_kernel(input, structure, output) elif center_is_true and not brute_force: raise NotImplementedError( 'only brute_force iteration has been implemented' ) else: if cupy.shares_memory(output, input, 'MAY_SHARE_BOUNDS'): raise ValueError('output and input may not overlap in memory') tmp_in = cupy.empty_like(input, dtype=output.dtype) tmp_out = output if iterations >= 1 and not iterations & 1: tmp_in, tmp_out = tmp_out, tmp_in if masked: tmp_out = erode_kernel(input, structure, mask, tmp_out) else: tmp_out = erode_kernel(input, structure, tmp_out) # TODO: kernel doesn't return the changed status, so determine it here changed = not (input == tmp_out).all() # synchronize! ii = 1 while ii < iterations or ((iterations < 1) and changed): tmp_in, tmp_out = tmp_out, tmp_in if masked: tmp_out = erode_kernel(tmp_in, structure, mask, tmp_out) else: tmp_out = erode_kernel(tmp_in, structure, tmp_out) changed = not (tmp_in == tmp_out).all() ii += 1 if not changed and (not ii & 1): # synchronize! # can exit early if nothing changed # (only do this after even number of tmp_in/out swaps) break output = tmp_out if temp_needed: _core.elementwise_copy(output, temp) output = temp return output
def fourier_ellipsoid(input, size, n=-1, axis=-1, output=None): """Multidimensional ellipsoid Fourier filter. The array is multiplied with the fourier transform of a ellipsoid of given sizes. Args: input (cupy.ndarray): The input array. size (float or sequence of float): The size of the box used for filtering. If a float, `size` is the same for all axes. If a sequence, `size` has to contain one value for each axis. n (int, optional): If `n` is negative (default), then the input is assumed to be the result of a complex fft. If `n` is larger than or equal to zero, the input is assumed to be the result of a real fft, and `n` gives the length of the array before transformation along the real transform direction. axis (int, optional): The axis of the real transform (only used when ``n > -1``). output (cupy.ndarray, optional): If given, the result of shifting the input is placed in this array. Returns: output (cupy.ndarray): The filtered output. """ ndim = input.ndim if ndim == 1: return fourier_uniform(input, size, n, axis, output) if ndim > 3: # Note: SciPy currently does not do any filtering on >=4d inputs, but # does not warn about this! raise NotImplementedError('Only 1d, 2d and 3d inputs are supported') output = _get_output_fourier(output, input) axis = internal._normalize_axis_index(axis, ndim) sizes = _util._fix_sequence_arg(size, ndim, 'size') _core.elementwise_copy(input, output) # compute the distance from the origin for all samples in Fourier space distance = 0 for ax, (size, ax_size) in enumerate(zip(sizes, output.shape)): # compute the frequency grid in Hz if ax == axis and n > 0: arr = cupy.arange(ax_size, dtype=output.real.dtype) arr *= numpy.pi * size / n else: arr = cupy.fft.fftfreq(ax_size) arr *= numpy.pi * size arr = arr.astype(output.real.dtype, copy=False) arr *= arr arr = _reshape_nd(arr, ndim=ndim, axis=ax) distance = distance + arr cupy.sqrt(distance, out=distance) if ndim == 2: special.j1(distance, out=output) output *= 2 output /= distance elif ndim == 3: cupy.sin(distance, out=output) output -= distance * cupy.cos(distance) output *= 3 output /= distance**3 output[(0, ) * ndim] = 1.0 # avoid NaN in corner at frequency=0 location output *= input return output