From 2898db6d1e35bd23c7d6a0f590e3995b004b3ddb Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Sat, 7 Oct 2017 01:15:19 +0200 Subject: [PATCH 01/38] ENH: add cupy tensor space --- odl/space/__init__.py | 3 + odl/space/cupy_tensors.py | 2248 +++++++++++++++++++++++++++++++++++++ odl/space/entry_points.py | 3 + 3 files changed, 2254 insertions(+) create mode 100644 odl/space/cupy_tensors.py diff --git a/odl/space/__init__.py b/odl/space/__init__.py index 36e99548a05..4b317575d43 100644 --- a/odl/space/__init__.py +++ b/odl/space/__init__.py @@ -19,6 +19,9 @@ from .npy_tensors import * __all__ += npy_tensors.__all__ +from .cupy_tensors import * +__all__ += cupy_tensors.__all__ + from .pspace import * __all__ += pspace.__all__ diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py new file mode 100644 index 00000000000..27480cd4543 --- /dev/null +++ b/odl/space/cupy_tensors.py @@ -0,0 +1,2248 @@ +# Copyright 2014-2017 The ODL contributors +# +# This file is part of ODL. +# +# This Source Code Form is subject to the terms of the Mozilla Public License, +# v. 2.0. If a copy of the MPL was not distributed with this file, You can +# obtain one at https://mozilla.org/MPL/2.0/. + +"""Implementation of tensor spaces using ``pygpu``.""" + +from __future__ import print_function, division, absolute_import +import numpy as np + +from odl.set import RealNumbers +from odl.space.base_tensors import TensorSpace, Tensor +from odl.space.weighting import ( + Weighting, ArrayWeighting, ConstWeighting, + CustomInner, CustomNorm, CustomDist) +from odl.util import dtype_str, is_floating_dtype, signature_string + +try: + import cupy +except ImportError: + CUPY_AVAILABLE = False +else: + from pkg_resources import parse_version + if parse_version(cupy.__version__) < parse_version('2.0.0rc1'): + raise ImportError('cupy <2.0.0rc1 not supported') + CUPY_AVAILABLE = True + + +__all__ = ('CupyTensorSpace',) + + +# --- Space method implementations --- # + + +lico = cupy.ElementwiseKernel(in_params='T a, T x, T b, T y', + out_params='T z', + operation='z = a * x + b * y;', + name='lico') + + +def fallback_scal(a, x): + x *= a + return x + + +def fallback_axpy(a, x, y): + return lico(a, x, 1, y, y) + + +def _flat_inc(arr): + """Compute the flat element stride for cuBLAS if possible, else raise.""" + flat_inc = min(arr.strides) // arr.itemsize + stride = min(arr.strides) + for n, s in zip(sorted(arr.shape)[:-1], sorted(arr.strides)[1:]): + next_stride = stride * n + if s != next_stride: + raise ValueError + return flat_inc + + +def _cublas_func(name, dtype): + """Return the specified cupy.cuda.cublas function for a given dtype. + + Parameters + ---------- + name : str + Raw function name without prefix, e.g., ``'axpy'``. + dtype : + Numpy dtype specifier for which the cuBLAS function should be + used. Must be either single or double precision float. + + Raises + ------ + ValueError : + If the data type is not supported by cuBLAS. + """ + if np.dtype(dtype) == 'float32': + prefix = 's' + elif np.dtype(dtype) == 'float64': + prefix = 'd' + else: + raise ValueError('dtype {!r} not supported by cuBLAS'.format(dtype)) + + return getattr(cupy.cuda.cublas, prefix + name) + + +def _get_scal_axpy(x1, x2): + """Return implementations of scal and axpy suitable for the inputs.""" + try: + incx1 = _flat_inc(x1.data) + incx2 = _flat_inc(x2.data) + except ValueError: + use_cublas = False + else: + use_cublas = True + + if use_cublas: + try: + scal_cublas = _cublas_func('scal', x1.dtype) + except (ValueError, AttributeError): + scal = fallback_scal + else: + def scal(a, x): + with cupy.cuda.Device(x1.device) as dev: + return scal_cublas( + dev.cublas_handle, x.data.size, a, x.data.ptr, incx1) + + try: + axpy_cublas = _cublas_func('axpy', x1.dtype) + except (ValueError, AttributeError): + axpy = fallback_axpy + else: + def axpy(a, x, y): + with cupy.cuda.Device(x1.device) as dev: + return axpy_cublas( + dev.cublas_handle, x.data.size, a, + x.data.ptr, incx1, y.data.ptr, incx2) + else: + scal = fallback_scal + axpy = fallback_axpy + + return scal, axpy + + +def _lincomb_impl(a, x1, b, x2, out): + """Linear combination implementation, assuming types have been checked. + + This implementation is a highly optimized, considering all special + cases of array alignment and special scalar values 0 and 1 separately. + """ + scal, axpy = _get_scal_axpy(x1, x2) + + if a == 0 and b == 0: + # out <- 0 + out.data.fill(0) + + elif a == 0: + # Compute out <- b * x2 + if out is x2: + # out <- b * out + if b == 1: + pass + else: + scal(b, out.data) + else: + # out <- b * x2 + if b == 1: + out.data[:] = x2.data + else: + cupy.multiply(b, x2.data, out=out.data) + + elif b == 0: + # Compute out <- a * x1 + if out is x1: + # out <- a * out + if a == 1: + pass + else: + scal(a, out.data) + else: + # out <- a * x1 + if a == 1: + out.data[:] = x1.data + else: + cupy.multiply(a, x1.data, out=out.data) + + else: + # Compute out <- a * x1 + b * x2 + # Optimize a number of alignment options. We know that a and b + # are nonzero. + if out is x1 and out is x2: + # out <-- (a + b) * out + if a + b == 0: + out.data.fill(0) + elif a + b == 1: + pass + else: + scal(a + b, out.data) + elif out is x1 and a == 1: + # out <-- out + b * x2 + axpy(b, x2.data, out.data) + elif out is x2 and b == 1: + # out <-- a * x1 + out + axpy(a, x1.data, out.data) + else: + # out <-- a * x1 + b * x2 + # No optimization for other cases of a and b; alignment doesn't + # matter anymore. + lico(a, x1.data, b, x2.data, out.data) + + +# --- Space and element classes --- # + + +class CupyTensorSpace(TensorSpace): + + """Tensor space implemented with GPU arrays. + + This space implements tensors of arbitrary rank over a `Field` ``F``, + which is either the real or complex numbers. + + Its elements are represented as instances of the + `CupyTensor` class. + """ + + def __init__(self, shape, dtype='float64', device=None, **kwargs): + """Initialize a new instance. + + Parameters + ---------- + shape : sequence of non-negative ints + Number entries per dimension. + dtype : + Data type for each tuple entry. Can be provided in any + way the `numpy.dtype` function understands, e.g., + as built-in type, as one of NumPy's internal datatype + objects or as string. + See `available_dtypes` for the list of supported data types. + device : int, optional + ID of the GPU device where elements should be created. + For ``None``, the default device is chosen, which usually + has ID 0. + weighting : optional + Use weighted inner product, norm, and dist. The following + types are supported: + + `Weighting`: Use this weighting as-is. + Compatibility with this space's elements is not checked + during init. + + float: Weighting by a constant + + array-like: Pointwise weighting by an array of the same + `shape` as the space. + + sequence of 1D array-likes: Per-axis (tensor product) weighting + using broadcasting multiplication in each axis. ``None`` + entries cause the corresponding axis to be skipped. + + This option cannot be combined with ``dist``, + ``norm`` or ``inner``. + + Default: no weighting + + exponent : positive float, optional + Exponent of the norm. For values other than 2.0, no + inner product is defined. + + This option is ignored if ``dist``, ``norm`` or + ``inner`` is given. + + Default: 2.0 + + Other Parameters + ---------------- + dist : callable, optional + The distance function defining a metric on the space. + It must accept two `CupyTensor` arguments and + fulfill the following mathematical conditions for any + three vectors ``x, y, z``: + + - ``dist(x, y) >= 0`` + - ``dist(x, y) = 0`` if and only if ``x = y`` + - ``dist(x, y) = dist(y, x)`` + - ``dist(x, y) <= dist(x, z) + dist(z, y)`` + + This option cannot be combined with ``weight``, + ``norm`` or ``inner``. + + norm : callable, optional + The norm implementation. It must accept an + `CupyTensor` argument, return a float and satisfy the + following conditions for all vectors ``x, y`` and scalars + ``s``: + + - ``||x|| >= 0`` + - ``||x|| = 0`` if and only if ``x = 0`` + - ``||s * x|| = |s| * ||x||`` + - ``||x + y|| <= ||x|| + ||y||`` + + By default, ``norm(x)`` is calculated as ``inner(x, x)``. + + This option cannot be combined with ``weight``, + ``dist`` or ``inner``. + + inner : callable, optional + The inner product implementation. It must accept two + `CupyTensor` arguments, return a element from + the field of the space (real or complex number) and + satisfy the following conditions for all vectors + ``x, y, z`` and scalars ``s``: + + - `` = conj()`` + - `` = s * + `` + - `` = 0`` if and only if ``x = 0`` + + This option cannot be combined with ``weight``, + ``dist`` or ``norm``. + + kwargs : + Further keyword arguments are passed to the weighting + classes. + + Examples + -------- + Initialization with the class constructor: + + >>> space = CupyTensorSpace(3, 'float') + >>> space + rn(3, impl='cupy') + >>> space.shape + (3,) + >>> space.dtype + dtype('float64') + + A more convenient way is to use the factory functions with the + ``impl='cupy'`` option: + + >>> space = odl.rn(3, impl='cupy', weighting=[1, 2, 3]) + >>> space + rn(3, impl='cupy', weighting=[1, 2, 3]) + >>> space = odl.tensor_space((2, 3), impl='cupy', dtype=int) + >>> space + tensor_space((2, 3), 'int', impl='cupy') + """ + super(CupyTensorSpace, self).__init__(shape, dtype) + if self.dtype.char not in self.available_dtypes(): + raise ValueError('`dtype` {!r} not supported'.format(dtype)) + + if device is None: + self.__device = cupy.cuda.get_device_id() + else: + self.__device = int(device) + + dist = kwargs.pop('dist', None) + norm = kwargs.pop('norm', None) + inner = kwargs.pop('inner', None) + weighting = kwargs.pop('weighting', None) + exponent = kwargs.pop('exponent', 2.0) + + # Check validity of option combination (3 or 4 out of 4 must be None) + if sum(x is None for x in (dist, norm, inner, weighting)) < 3: + raise ValueError('invalid combination of options `weighting`, ' + '`dist`, `norm` and `inner`') + if any(x is not None for x in (dist, norm, inner)) and exponent != 2.0: + raise ValueError('`exponent` cannot be used together with ' + '`dist`, `norm` and `inner`') + + # Set the weighting + if weighting is not None: + if isinstance(weighting, Weighting): + if weighting.impl != 'cupy': + raise ValueError("`weighting.impl` must be 'cupy', " + '`got {!r}'.format(weighting.impl)) + if weighting.exponent != exponent: + raise ValueError('`weighting.exponent` conflicts with ' + '`exponent`: {} != {}' + ''.format(weighting.exponent, exponent)) + self.__weighting = weighting + else: + self.__weighting = _weighting(weighting, exponent) + + # Check (afterwards) that the weighting input was sane + if isinstance(self.weighting, CupyTensorSpaceArrayWeighting): + if not np.can_cast(self.weighting.array.dtype, self.dtype): + raise ValueError( + 'cannot cast from `weighting` data type {} to ' + 'the space `dtype` {}' + ''.format(dtype_str(self.weighting.array.dtype), + dtype_str(self.dtype))) + if self.weighting.array.shape != self.shape: + raise ValueError('array-like weights must have same ' + 'shape {} as this space, got {}' + ''.format(self.shape, + self.weighting.array.shape)) + + elif dist is not None: + self.__weighting = CupyTensorSpaceCustomDist(dist) + elif norm is not None: + self.__weighting = CupyTensorSpaceCustomNorm(norm) + elif inner is not None: + self.__weighting = CupyTensorSpaceCustomInner(inner) + else: # all None -> no weighing + self.__weighting = CupyTensorSpaceConstWeighting(1.0, exponent) + + @property + def device(self): + """The GPU device ID of this tensor space.""" + return self.__device + + @property + def impl(self): + """Implementation back-end of this space: ``'cupy'``.""" + return 'cupy' + + @property + def default_order(self): + """Default storage order for new elements in this space: ``'C'``.""" + return 'C' + + @property + def weighting(self): + """This space's weighting scheme.""" + return self.__weighting + + @property + def is_weighted(self): + """Return ``True`` if the space is not weighted by constant 1.0.""" + return not ( + isinstance(self.weighting, CupyTensorSpaceConstWeighting) and + self.weighting.const == 1.0) + + @property + def exponent(self): + """Exponent of the norm and distance.""" + return self.weighting.exponent + + def element(self, inp=None, order=None): + """Create a new element. + + Parameters + ---------- + inp : `array-like`, optional + Input used to initialize the new element. + + If ``inp`` is `None`, an empty element is created with no + guarantee of its state (memory allocation only). + The new element will use ``order`` as storage order if + provided, otherwise `default_order`. + + Otherwise, a copy is avoided whenever possible. This requires + correct `shape` and `dtype`, and if ``order`` is provided, + also contiguousness in that ordering. If any of these + conditions is not met, a copy is made. + + order : {None, 'C', 'F'}, optional + Storage order of the returned element. For ``'C'`` and ``'F'``, + contiguous memory in the respective ordering is enforced. + The default ``None`` enforces no contiguousness. + + Returns + ------- + element : `CupyTensor` + The new element created (from ``inp``). + + Notes + ----- + This method preserves "array views" of correct size and type, + see the examples below. + + Examples + -------- + >>> space = odl.rn((2, 3), impl='cupy') + + Create an empty element: + + >>> empty = space.element() + >>> empty.shape + (2, 3) + + Initialization during creation: + + >>> x = space.element([[1, 2, 3], + ... [4, 5, 6]]) + >>> x + rn((2, 3), impl='cupy').element( + [[ 1., 2., 3.], + [ 4., 5., 6.]] + ) + """ + if order is None: + order_in = order + else: + order, order_in = str(order).upper(), order + + if order is not None and order not in ('C', 'F'): + raise ValueError("`order` {!r} not understood".format(order_in)) + + with cupy.cuda.Device(self.device): + if inp is None: + if order is None: + order = self.default_order + arr = cupy.empty(self.shape, dtype=self.dtype, order=order) + + else: + if inp in self and order is None: + # Short-circuit for space elements and no enforced ordering + return inp + + if hasattr(inp, 'shape') and inp.shape != self.shape: + raise ValueError('`inp` must have shape {}, got shape {}' + ''.format(self.shape, inp.shape)) + + if isinstance(inp, cupy.ndarray): + # Workaround for https://github.com/cupy/cupy/issues/590 + # TODO: remove when solved + if (inp.dtype == self.dtype and + inp.device.id == self.device): + arr = inp + else: + arr = inp.astype(self.dtype) + else: + arr = cupy.array(inp, copy=False, dtype=self.dtype, + ndmin=self.ndim, order=order) + + # If the result has a 0 stride, make a copy since it would + # produce all kinds of nasty problems. This happens for e.g. + # results of `broadcast_to()`. + if 0 in arr.strides: + arr = arr.copy() + + return self.element_type(self, arr) + + def zero(self): + """Create a tensor filled with zeros. + + Examples + -------- + >>> space = odl.rn(3, impl='cupy') + >>> x = space.zero() + >>> x + rn(3, impl='cupy').element([ 0., 0., 0.]) + """ + with cupy.cuda.Device(self.device): + arr = cupy.zeros(self.shape, dtype=self.dtype) + return self.element(arr) + + def one(self): + """Create a tensor filled with ones. + + Examples + -------- + >>> space = odl.rn(3, impl='cupy') + >>> x = space.one() + >>> x + rn(3, impl='cupy').element([ 1., 1., 1.]) + """ + with cupy.cuda.Device(self.device): + arr = cupy.ones(self.shape, dtype=self.dtype) + return self.element(arr) + + def __eq__(self, other): + """Return ``self == other``. + + Returns + ------- + equals : bool + ``True`` if ``other`` is an instance of this space's type + with the same `shape`, `dtype`, `device` and + `weighting`, ``False`` otherwise. + + Examples + -------- + >>> space = odl.rn(2, impl='cupy') + >>> same_space = odl.rn(2, exponent=2, impl='cupy') + >>> same_space == space + True + + Different `shape`, `exponent`, `dtype` or `impl` + all result in different spaces: + + >>> diff_space = odl.rn((2, 3), impl='cupy') + >>> diff_space == space + False + >>> diff_space = odl.rn(2, exponent=1, impl='cupy') + >>> diff_space == space + False + >>> diff_space = odl.rn(2, dtype='float32', impl='cupy') + >>> diff_space == space + False + >>> diff_space = odl.rn(2, impl='numpy') + >>> diff_space == space + False + >>> space == object + False + + A `CupyTensorSpace` with the same properties is considered + equal: + + >>> same_space = odl.CupyTensorSpace(2, dtype='float64') + >>> same_space == space + True + """ + return (super().__eq__(other) and + self.device == other.device and + self.weighting == other.weighting) + + def __hash__(self): + """Return ``hash(self)``.""" + return hash((super().__hash__(), self.device, self.weighting)) + + def _lincomb(self, a, x1, b, x2, out): + """Linear combination of ``x1`` and ``x2``. + + Calculate ``out = a*x1 + b*x2`` using optimized BLAS + routines if possible. + + Parameters + ---------- + a, b : `TensorSpace.field` elements + Scalars to multiply ``x1`` and ``x2`` with. + x1, x2 : `CupyTensor` + Summands in the linear combination. + out : `CupyTensor` + Tensor to which the result is written. + + Returns + ------- + None + + Examples + -------- + >>> r3 = odl.rn(3, impl='cupy') + >>> x = r3.element([1, 2, 3]) + >>> y = r3.element([4, 5, 6]) + >>> out = r3.element() + >>> result = r3.lincomb(2, x, -1, y, out) + >>> result + rn(3, impl='cupy').element([-2., -1., 0.]) + >>> result is out + True + """ + _lincomb_impl(a, x1, b, x2, out) + + def _dist(self, x1, x2): + """Calculate the distance between two tensors. + + Parameters + ---------- + x1, x2 : `CupyTensor` + Tensors whose mutual distance is calculated. + + Returns + ------- + dist : float + Distance between the tensors. + + Examples + -------- + The default case is the Euclidean distance: + + >>> r3 = odl.rn(3, impl='cupy') + >>> x = r3.element([1, 2, 3]) + >>> y = r3.element([4, 2, -1]) + >>> r3.dist(x, y) # 3^2 + 4^2 = 25 + 5.0 + + Taking a different exponent or a weighting is also possible + during space creation: + + >>> r3 = odl.rn(3, impl='cupy', exponent=1) + >>> x = r3.element([1, 2, 3]) + >>> y = r3.element([4, 2, -1]) + >>> r3.dist(x, y) # 3 + 4 = 7 + 7.0 + + >>> r3 = odl.rn(3, impl='cupy', weighting=2, exponent=1) + >>> x = r3.element([1, 2, 3]) + >>> y = r3.element([4, 2, -1]) + >>> r3.dist(x, y) # 2*3 + 2*4 = 14 + 14.0 + """ + return self.weighting.dist(x1, x2) + + def _norm(self, x): + """Calculate the norm of a tensor. + + Parameters + ---------- + x : `CupyTensor` + The tensor whose norm is calculated. + + Returns + ------- + norm : float + Norm of the tensor. + + Examples + -------- + The default case is the Euclidean norm: + + >>> r3 = odl.rn(3, impl='cupy') + >>> x = r3.element([3, 4, 0]) + >>> r3.norm(x) # 3^2 + 4^2 = 25 + 5.0 + + Taking a different exponent or a weighting is also possible + during space creation: + + >>> r3 = odl.rn(3, impl='cupy', exponent=1) + >>> x = r3.element([3, 4, 0]) + >>> r3.norm(x) # 3 + 4 = 7 + 7.0 + + >>> r3 = odl.rn(3, impl='cupy', weighting=2, exponent=1) + >>> x = r3.element([3, 4, 0]) + >>> r3.norm(x) # 2*3 + 2*4 = 14 + 14.0 + """ + return self.weighting.norm(x) + + def _inner(self, x1, x2): + """Raw inner product of two tensors. + + Parameters + ---------- + x1, x2 : `CupyTensor` + The tensors whose inner product is calculated. + + Returns + ------- + inner : `field` element + Inner product of the tensors. + + Examples + -------- + The default case is the dot product: + + >>> r3 = odl.rn(3, impl='cupy') + >>> x = r3.element([1, 2, 3]) + >>> y = r3.element([-1, 0, 1]) + >>> r3.inner(x, y) # 1*(-1) + 2*0 + 3*1 = 2 + 2.0 + + Taking a different weighting is also possible during space + creation: + + >>> r3 = odl.rn(3, impl='cupy', weighting=2) + >>> x = r3.element([1, 2, 3]) + >>> y = r3.element([-1, 0, 1]) + >>> r3.inner(x, y) # 2 * 1*(-1) + 2 * 2*0 + 2 * 3*1 = 4 + 4.0 + """ + return self.weighting.inner(x1, x2) + + def _multiply(self, x1, x2, out): + """Entry-wise product of two tensors, assigned to out. + + Parameters + ---------- + x1, x2 : `CupyTensor` + Factors in the product. + out : `CupyTensor` + Tensor to which the result is written. + + Examples + -------- + Out-of-place evaluation: + + >>> r3 = odl.rn(3, impl='cupy') + >>> x = r3.element([1, 2, 3]) + >>> y = r3.element([-1, 0, 1]) + >>> r3.multiply(x, y) + rn(3, impl='cupy').element([-1., 0., 3.]) + + In-place: + + >>> out = r3.element() + >>> result = r3.multiply(x, y, out=out) + >>> result + rn(3, impl='cupy').element([-1., 0., 3.]) + >>> result is out + True + """ + x1.ufuncs.multiply(x2, out=out) + + def _divide(self, x1, x2, out): + """Entry-wise division of two tensors, assigned to out. + + Parameters + ---------- + x1, x2 : `CupyTensor` + Dividend and divisor in the quotient. + out : `CupyTensor` + Tensor to which the result is written. + + Examples + -------- + Out-of-place evaluation: + + >>> r3 = odl.rn(3, impl='cupy') + >>> x = r3.element([1, 2, 3]) + >>> y = r3.element([-1, 2, 1]) + >>> r3.divide(x, y) + rn(3, impl='cupy').element([-1., 1., 3.]) + + In-place: + + >>> out = r3.element() + >>> result = r3.divide(x, y, out=out) + >>> result + rn(3, impl='cupy').element([-1., 1., 3.]) + >>> result is out + True + """ + x1.ufuncs.divide(x2, out=out) + + def __repr__(self): + """Return ``repr(self)``.""" + if self.ndim == 1: + posargs = [self.size] + else: + posargs = [self.shape] + + if self.is_real: + constructor_name = 'rn' + elif self.is_complex: + constructor_name = 'cn' + else: + constructor_name = 'tensor_space' + + if (constructor_name == 'tensor_space' or + (not self.is_real and not self.is_complex) or + self.dtype != self.default_dtype(self.field)): + posargs.append(dtype_str(self.dtype)) + + optargs = [('impl', self.impl, 'numpy'), # for the helper functions + ('device', self.device, cupy.cuda.get_device_id())] + inner_str = signature_string(posargs, optargs) + weight_str = self.weighting.repr_part + if weight_str: + inner_str += ', ' + weight_str + + return '{}({})'.format(constructor_name, inner_str) + + @property + def element_type(self): + """`CupyTensor`""" + return CupyTensor + + @staticmethod + def available_dtypes(): + """Return the data types available for this space.""" + dtypes = (np.sctypes['float'] + + np.sctypes['complex'] + + np.sctypes['int'] + + np.sctypes['uint'] + + [bool]) + dtypes.remove(np.float128) + dtypes.remove(np.complex256) + return tuple(np.dtype(dtype) for dtype in dtypes) + + @staticmethod + def default_dtype(field=None): + """Return the default data type of this space type for a given field. + + Parameters + ---------- + field : `Field`, optional + Set of numbers to be represented by a data type. + Currently supported : `RealNumbers`, `ComplexNumbers`. + Default: `RealNumbers` + + Returns + ------- + dtype : `numpy.dtype` + Numpy data type specifier. The returned defaults are: + + ``RealNumbers()`` : ``np.dtype('float64')`` + + ``ComplexNumbers()`` : not supported + """ + if field is None or field == RealNumbers(): + return np.dtype('float64') + else: + raise ValueError('no default data type defined for field {}.' + ''.format(field)) + + +class CupyTensor(Tensor): + + """Representation of an `CupyTensorSpace` element.""" + + def __init__(self, space, data): + """Initialize a new instance.""" + super(CupyTensor, self).__init__(space) + self.__data = data + + @property + def data(self): + """Raw `cupy.core.core.ndarray` representing the data.""" + return self.__data + + @property + def ndim(self): + """Number of axes (=dimensions) of this tensor.""" + return self.space.ndim + + @property + def device(self): + """The GPU device on which this tensor lies.""" + return self.space.device + + def asarray(self, out=None): + """Extract the data of this element as a `numpy.ndarray`. + + Parameters + ---------- + out : `numpy.ndarray`, optional + Array to which the result should be written. + Has to be contiguous and of the correct data type. + + Returns + ------- + asarray : `numpy.ndarray` + Numpy array of the same `dtype` and `shape` this tensor. + If ``out`` was given, the returned object is a reference to it. + + Examples + -------- + By default, a new array is created: + + >>> r3 = odl.rn(3, impl='cupy') + >>> x = r3.element([1, 2, 3]) + >>> x.asarray() + array([ 1., 2., 3.]) + >>> int_spc = odl.tensor_space(3, impl='cupy', dtype=int) + >>> x = int_spc.element([1, 2, 3]) + >>> x.asarray() + array([1, 2, 3]) + >>> tensors = odl.rn((2, 3), impl='cupy', dtype='float32') + >>> x = tensors.element([[1, 2, 3], + ... [4, 5, 6]]) + >>> x.asarray() + array([[ 1., 2., 3.], + [ 4., 5., 6.]], dtype=float32) + + Using the out parameter, the array can be filled in-place: + + >>> out = np.empty((2, 3), dtype='float32') + >>> result = x.asarray(out=out) + >>> out + array([[ 1., 2., 3.], + [ 4., 5., 6.]], dtype=float32) + >>> result is out + True + """ + if out is None: + return cupy.asnumpy(self.data) + else: + if out.shape != self.shape: + raise ValueError('`out` must have shape {}, got shape {}' + ''.format(self.shape, out.shape)) + if out.dtype != self.dtype: + raise ValueError('`out` must have dtype {}, got dtype {}' + ''.format(self.dtype, out.dtype)) + self.data.data.copy_to_host( + out.ctypes.data_as(np.ctypeslib.ctypes.c_void_p), + self.size * self.itemsize) + return out + + @property + def data_ptr(self): + """A raw pointer to the data container. + + Examples + -------- + >>> r3 = odl.rn(3, impl='cupy') + >>> x = r3.one() + >>> x.data_ptr # doctest: +SKIP + 47259975936 + """ + return self.data.ptr + + def __eq__(self, other): + """Return ``self == other``. + + Parameters + ---------- + other : + Object to be compared with ``self``. + + Returns + ------- + equals : bool + ``True`` if all entries of ``other`` are equal to this + tensor's entries, ``False`` otherwise. + + Notes + ----- + The element-by-element comparison is performed on the CPU, + i.e. it involves data transfer to host memory, which is slow. + + Examples + -------- + >>> r3 = odl.rn(3, impl='cupy') + >>> x = r3.element([1, 2, 3]) + >>> same_x = r3.element([1, 2, 3]) + >>> y = r3.element([-1, -2, -3]) + >>> x == same_x + True + >>> x == y + False + + Space membership matters: + + >>> int_spc = odl.tensor_space(3, impl='cupy', dtype=int) + >>> x_int = int_spc.element([1, 2, 3]) + >>> x == x_int + False + """ + if other is self: + return True + elif other not in self.space: + return False + else: + return bool((self.data == other.data).all()) + + def copy(self): + """Create an identical (deep) copy of this tensor. + + Returns + ------- + copy : `pygpu._array.ndgpuarray` + A deep copy. + + Examples + -------- + >>> r3 = odl.rn(3, impl='cupy') + >>> x = r3.element([1, 2, 3]) + >>> y = x.copy() + >>> y + rn(3, impl='cupy').element([ 1., 2., 3.]) + >>> x == y + True + >>> x is y + False + """ + return self.space.element(self.data.copy()) + + def __getitem__(self, indices): + """Access values of this tensor. + + Parameters + ---------- + indices : index expression + The position(s) that should be accessed. + + Returns + ------- + values : scalar or `pygpu._array.ndgpuarray` + The value(s) at the index (indices). + + Examples + -------- + Indexing rules follow roughly the Numpy style, as far (or "fancy") + as supported: + + >>> r5 = odl.rn(5, impl='cupy') + >>> x = r5.element([1, 2, 3, 4, 5]) + >>> x[1:4] + rn(3, impl='cupy').element([ 2., 3., 4.]) + >>> x[::2] + rn(3, impl='cupy').element([ 1., 3., 5.]) + + The returned views are writable, so modificatons alter the + original array: + + >>> view = x[1:4] + >>> view[:] = -1 + >>> view + rn(3, impl='cupy').element([-1., -1., -1.]) + >>> x + rn(5, impl='cupy').element([ 1., -1., -1., -1., 5.]) + + Multi-indexing is also directly supported: + + >>> tensors = odl.rn((2, 3), impl='cupy') + >>> x = tensors.element([[1, 2, 3], + ... [4, 5, 6]]) + >>> x[1, 2] + 6.0 + >>> x[1] # row with index 1 + rn(3, impl='cupy').element([ 4., 5., 6.]) + >>> view = x[:, ::2] + >>> view + rn((2, 2), impl='cupy').element( + [[ 1., 3.], + [ 4., 6.]] + ) + >>> view[:] = [[0, 0], + ... [0, 0]] + >>> x + rn((2, 3), impl='cupy').element( + [[ 0., 2., 0.], + [ 0., 5., 0.]] + ) + """ + arr = self.data[indices] + if arr.shape == (): + if arr.dtype.kind == 'f': + return float(np.asarray(arr)) + elif arr.dtype.kind == 'c': + return complex(np.asarray(arr)) + elif arr.dtype.kind in ('u', 'i'): + return int(np.asarray(arr)) + else: + raise RuntimeError("no conversion for dtype {}" + "".format(arr.dtype)) + else: + space = type(self.space)(arr.shape, dtype=self.dtype, + device=self.device) + return space.element(arr) + + def __setitem__(self, indices, values): + """Set values of this tensor. + + Parameters + ---------- + indices : index expression + The position(s) that should be accessed. + values : scalar or `array-like` + The value(s) that are to be assigned. + + If ``indices`` is an int (1D) or a sequence of ints, + ``value`` must be scalar. + + Otherwise, ``value`` must be broadcastable to the shape of + the sliced view according to the Numpy broadcasting rules. + + Examples + -------- + In 1D, Values can be set with scalars or arrays that match the + shape of the slice: + + >>> r5 = odl.rn(5, impl='cupy') + >>> x = r5.element([1, 2, 3, 4, 5]) + >>> x[1:4] = 0 + >>> x + rn(5, impl='cupy').element([ 1., 0., 0., 0., 5.]) + >>> x[1:4] = [-1, 1, -1] + >>> x + rn(5, impl='cupy').element([ 1., -1., 1., -1., 5.]) + >>> y = r5.element([5, 5, 5, 8, 8]) + >>> x[:] = y + >>> x + rn(5, impl='cupy').element([ 5., 5., 5., 8., 8.]) + + In higher dimensions, broadcasting can be applied to assign + values: + + >>> tensors = odl.rn((2, 3), impl='cupy') + >>> x = tensors.element([[1, 2, 3], + ... [4, 5, 6]]) + >>> x[:] = [[6], [3]] # rhs mimics (2, 1) shape + >>> x + rn((2, 3), impl='cupy').element( + [[ 6., 6., 6.], + [ 3., 3., 3.]] + ) + + Be aware of unsafe casts and over-/underflows, there + will be warnings at maximum. + + >>> int_r3 = odl.tensor_space(3, impl='cupy', dtype='uint32') + >>> x = int_r3.element([1, 2, 3]) + >>> x[0] = -1 + >>> x[0] + 4294967295 + """ + if isinstance(values, CupyTensor): + self.data[indices] = values.data + elif np.isscalar(values): + self.data[indices] = values + else: + values = cupy.array(values, dtype=self.dtype, copy=False) + self.data[indices] = values + + def __int__(self): + """Return ``int(self)``. + + Returns + ------- + int : int + Integer representing this tensor. + + Raises + ------ + TypeError + If the tensor is of `size` != 1. + """ + if self.size != 1: + raise TypeError('only size 1 tensors can be converted to int') + return int(self[(0,) * self.ndim]) + + def __long__(self): + """Return ``long(self)``. + + The `long` method is only available in Python 2. + + Returns + ------- + long : `long` + Integer representing this tensor. + + Raises + ------ + TypeError + If the tensor is of `size` != 1. + """ + if self.size != 1: + raise TypeError('only size 1 tensors can be converted to long') + return long(self[(0,) * self.ndim]) + + def __float__(self): + """Return ``float(self)``. + + Returns + ------- + float : float + Floating point number representing this tensor. + + Raises + ------ + TypeError + If the tensor is of `size` != 1. + """ + if self.size != 1: + raise TypeError('only size 1 tensors can be converted to float') + return float(self[(0,) * self.ndim]) + + def __complex__(self): + """Return ``complex(self)``. + + Returns + ------- + complex : `complex` + Complex floating point number representing this tensor. + + Raises + ------ + TypeError + If the tensor is of `size` != 1. + """ + if self.size != 1: + raise TypeError('only size 1 tensors can be converted to complex') + return complex(self[(0,) * self.ndim]) + + def __str__(self): + """Return ``str(self)``.""" + return str(self.data) + + def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): + """Interface to Numpy's ufunc machinery. + + This method is called by Numpy version 1.13 and higher as a single + point for the ufunc dispatch logic. An object implementing + ``__array_ufunc__`` takes over control when a `numpy.ufunc` is + called on it, allowing it to use custom implementations and + output types. + + This includes handling of in-place arithmetic like + ``npy_array += custom_obj``. In this case, the custom object's + ``__array_ufunc__`` takes precedence over the baseline + `numpy.ndarray` implementation. It will be called with + ``npy_array`` as ``out`` argument, which ensures that the + returned object is a Numpy array. For this to work properly, + ``__array_ufunc__`` has to accept Numpy arrays as ``out`` arguments. + The same holds analogously for GPU arrays. + + See the `corresponding NEP`_ and the `interface documentation`_ + for further details. See also the `general documentation on + Numpy ufuncs`_. + + .. note:: + This implementation looks for native ufuncs in ``pygpu.ufuncs`` + and falls back to the basic implementation with Numpy arrays + in case no native ufunc is available. That fallback version + comes with significant overhead due to data copies between + host and device. + + .. note:: + When an ``out`` parameter is specified, and (one of) it has + type `numpy.ndarray`, the inputs are converted to Numpy + arrays, and the Numpy ufunc is invoked. + + .. note:: + When using operations that alter the shape (like ``reduce``), + or the data type (can be any of the methods), + the resulting array is wrapped in a space of the same + type as ``self.space``, propagating space properties like + `exponent` or `weighting` as closely as possible. + + Parameters + ---------- + ufunc : `numpy.ufunc` + Ufunc that should be called on ``self``. + method : str + Method on ``ufunc`` that should be called on ``self``. + Possible values: + + ``'__call__'``, ``'accumulate'``, ``'at'``, ``'outer'``, + ``'reduce'``, ``'reduceat'`` + + input1, ..., inputN : + Positional arguments to ``ufunc.method``. + kwargs : + Keyword arguments to ``ufunc.method``. + + Returns + ------- + ufunc_result : `CupyTensor`, `numpy.ndarray` or tuple + Result of the ufunc evaluation. If no ``out`` keyword argument + was given, the result is a `Tensor` or a tuple + of such, depending on the number of outputs of ``ufunc``. + If ``out`` was provided, the returned object or tuple entries + refer(s) to ``out``. + + Examples + -------- + We apply `numpy.add` to ODL tensors: + + >>> r3 = odl.rn(3, impl='cupy') + >>> x = r3.element([1, 2, 3]) + >>> y = r3.element([-1, -2, -3]) + >>> x.__array_ufunc__(np.add, '__call__', x, y) + rn(3, impl='cupy').element([ 0., 0., 0.]) + >>> np.add(x, y) # same mechanism for Numpy >= 1.13 + rn(3, impl='cupy').element([ 0., 0., 0.]) + + As ``out``, a Numpy array or an ODL tensor can be given (wrapped + in a sequence): + + >>> out = r3.element() + >>> res = x.__array_ufunc__(np.add, '__call__', x, y, out=(out,)) + >>> out + rn(3, impl='cupy').element([ 0., 0., 0.]) + >>> res is out + True + >>> out_arr = np.empty(3) + >>> res = x.__array_ufunc__(np.add, '__call__', x, y, out=(out_arr,)) + >>> out_arr + array([ 0., 0., 0.]) + >>> res is out_arr + True + + With multiple dimensions: + + >>> r23 = odl.rn((2, 3), impl='cupy') + >>> x = y = r23.one() + >>> x.__array_ufunc__(np.add, '__call__', x, y) + rn((2, 3), impl='cupy').element( + [[ 2., 2., 2.], + [ 2., 2., 2.]] + ) + + The ``ufunc.accumulate`` method retains the original `shape` and + `dtype`. The latter can be changed with the ``dtype`` parameter: + + >>> x = r3.element([1, 2, 3]) + >>> x.__array_ufunc__(np.add, 'accumulate', x) + rn(3, impl='cupy').element([ 1., 3., 6.]) + >>> np.add.accumulate(x) # same mechanism for Numpy >= 1.13 + rn(3, impl='cupy').element([ 1., 3., 6.]) + + For multi-dimensional tensors, an optional ``axis`` parameter + can be provided: + + >>> z = r23.one() + >>> z.__array_ufunc__(np.add, 'accumulate', z, axis=1) + rn((2, 3), impl='cupy').element( + [[ 1., 2., 3.], + [ 1., 2., 3.]] + ) + + The ``ufunc.at`` method operates in-place. Here we add the second + operand ``[5, 10]`` to ``x`` at indices ``[0, 2]``: + + >>> x = r3.element([1, 2, 3]) + >>> x.__array_ufunc__(np.add, 'at', x, [0, 2], [5, 10]) + >>> x + rn(3, impl='cupy').element([ 6., 2., 13.]) + + For outer-product-type operations, i.e., operations where the result + shape is the sum of the individual shapes, the ``ufunc.outer`` + method can be used: + + >>> x = odl.rn(2, impl='cupy').element([0, 3]) + >>> y = odl.rn(3, impl='cupy').element([1, 2, 3]) + >>> x.__array_ufunc__(np.add, 'outer', x, y) + rn((2, 3), impl='cupy').element( + [[ 1., 2., 3.], + [ 4., 5., 6.]] + ) + >>> y.__array_ufunc__(np.add, 'outer', y, x) + rn((3, 2), impl='cupy').element( + [[ 1., 4.], + [ 2., 5.], + [ 3., 6.]] + ) + + Using ``ufunc.reduce`` produces a scalar, which can be avoided with + ``keepdims=True``: + + >>> x = r3.element([1, 2, 3]) + >>> x.__array_ufunc__(np.add, 'reduce', x) + 6.0 + >>> x.__array_ufunc__(np.add, 'reduce', x, keepdims=True) + rn(1, impl='cupy').element([ 6.]) + + In multiple dimensions, ``axis`` can be provided for reduction over + selected axes: + + >>> z = r23.element([[1, 2, 3], + ... [4, 5, 6]]) + >>> z.__array_ufunc__(np.add, 'reduce', z, axis=1) + rn(2, impl='cupy').element([ 6., 15.]) + + Finally, ``add.reduceat`` is a combination of ``reduce`` and + ``at`` with rather flexible and complex semantics (see the + `reduceat documentation`_ for details): + + >>> x = r3.element([1, 2, 3]) + >>> x.__array_ufunc__(np.add, 'reduceat', x, [0, 1]) + rn(2, impl='cupy').element([ 1., 5.]) + + References + ---------- + .. _corresponding NEP: + https://github.com/numpy/numpy/blob/master/doc/neps/\ +ufunc-overrides.rst + + .. _interface documentation: + https://github.com/charris/numpy/blob/master/doc/source/reference/\ +arrays.classes.rst#special-attributes-and-methods + + .. _general documentation on Numpy ufuncs: + https://docs.scipy.org/doc/numpy/reference/ufuncs.html + + .. _reduceat documentation: + https://docs.scipy.org/doc/numpy/reference/generated/\ +numpy.ufunc.reduceat.html + """ + # --- Process `out` --- # + + # Unwrap out if provided. The output parameters are all wrapped + # in one tuple, even if there is only one. + out_tuple = kwargs.pop('out', ()) + + # Check number of `out` args, depending on `method` + if method == '__call__' and len(out_tuple) not in (0, ufunc.nout): + raise ValueError( + "need 0 or {} `out` arguments for `method='__call__'`, " + 'got {}'.format(ufunc.nout, len(out_tuple))) + elif method != '__call__' and len(out_tuple) not in (0, 1): + raise ValueError( + "need 0 or 1 `out` arguments for `method={!r}`, " + 'got {}'.format(method, len(out_tuple))) + + # We allow our own tensors, the data container type and + # `numpy.ndarray` objects as `out` (see docs for reason for the + # latter) + valid_types = (type(self), type(self.data), np.ndarray) + if not all(isinstance(o, valid_types) or o is None + for o in out_tuple): + return NotImplemented + + # Determine native ufunc vs. Numpy ufunc + if any(isinstance(o, np.ndarray) for o in out_tuple): + native_ufunc = None + use_native = False + else: + native_ufunc = getattr(cupy, ufunc.__name__, None) + use_native = (native_ufunc is not None and + hasattr(native_ufunc, method)) + + # Assign to `out` or `out1` and `out2`, respectively, unwrapping the + # data container + out = out1 = out2 = None + if len(out_tuple) == 1: + if isinstance(out_tuple[0], type(self)): + out = out_tuple[0].data + else: + out = out_tuple[0] + elif len(out_tuple) == 2: + if isinstance(out_tuple[0], type(self)): + out1 = out_tuple[0].data + else: + out1 = out_tuple[0] + if isinstance(out_tuple[1], type(self)): + out1 = out_tuple[1].data + else: + out1 = out_tuple[1] + + # --- Process `inputs` --- # + + # Pull out the data container of the inputs if necessary + inputs = tuple( + inp.data if isinstance(inp, type(self)) else inp + for inp in inputs) + + # For native ufuncs, we turn non-scalar inputs into cupy arrays, + # as a workaround for https://github.com/cupy/cupy/issues/594 + # TODO: remove code when the upstream issue is fixed + if use_native: + inputs, orig_inputs = [], inputs + for inp in orig_inputs: + if (isinstance(inp, cupy.ndarray) or + np.isscalar(inp) or + inp is None): + inputs.append(inp) + else: + inputs.append(cupy.array(inp)) + + # --- Get some parameters for later --- # + + # Arguments for space constructors + exponent = self.space.exponent + weighting = self.space.weighting + + # --- Evaluate ufunc --- # + + if method == '__call__': + if ufunc.nout == 1: + if use_native: + kwargs['out'] = out # No tuple packing for cupy + res = native_ufunc(*inputs, **kwargs) + else: + kwargs['out'] = (out,) + # Everything is cast to Numpy arrays by the parent method; + # the result can be a Numpy array or a tensor + res = super(CupyTensor, self).__array_ufunc__( + ufunc, '__call__', *inputs, **kwargs) + + # Wrap result if necessary (lazily) + if out is None: + if is_floating_dtype(res.dtype): + # Weighting contains exponent + spc_kwargs = {'weighting': weighting} + else: + # No `exponent` or `weighting` applicable + spc_kwargs = {} + out_space = type(self.space)( + self.shape, res.dtype, self.device, **spc_kwargs) + return out_space.element(res) + else: + # `out` may be the unwrapped version, return the original + return out_tuple[0] + + elif ufunc.nout == 2: + kwargs['out'] = (out1, out2) + if use_native: + res1, res2 = native_ufunc(*inputs, **kwargs) + else: + # Everything is cast to Numpy arrays by the parent method; + # the results can be Numpy arrays or tensors + res1, res2 = super(CupyTensor, self).__array_ufunc__( + ufunc, '__call__', *inputs, **kwargs) + + # Wrap results if necessary (lazily) + # We don't use exponents or weightings since we don't know + # how to map them to the spaces + if out1 is None: + res_space = type(self.space)( + self.shape, res1.dtype, self.device) + result1 = res_space.element(res1) + else: + result1 = out_tuple[0] + + if out2 is None: + res_space = type(self.space)( + self.shape, res2.dtype, self.device) + result2 = res_space.element(res2) + else: + result2 = out_tuple[1] + + return result1, result2 + + else: + raise NotImplementedError('nout = {} not supported' + ''.format(ufunc.nout)) + + elif method == 'at': + native_method = getattr(native_ufunc, 'at', None) + use_native = (use_native and native_method is not None) + + def eval_at_via_npy(*inputs, **kwargs): + import ctypes + cupy_arr = inputs[0] + npy_arr = np.asarray(cupy_arr) + new_inputs = (npy_arr,) + inputs[1:] + super(CupyTensor, self).__array_ufunc__( + ufunc, method, *new_inputs, **kwargs) + # Workaround for https://github.com/cupy/cupy/issues/593 + # TODO: use cupy_arr[:] = npy_arr when it's fixed and not + # slower + cupy_arr.data.copy_from_host( + npy_arr.ctypes.data_as(ctypes.c_void_p), npy_arr.nbytes) + + if use_native: + # Native method could exist but raise `NotImplementedError` + # or return `NotImplemented`, falling back to Numpy case + # then, too + try: + res = native_method(*inputs, **kwargs) + except NotImplementedError: + eval_at_via_npy(*inputs, **kwargs) + else: + if res is NotImplemented: + eval_at_via_npy(*inputs, **kwargs) + else: + eval_at_via_npy(*inputs, **kwargs) + + else: # method != '__call__' + kwargs['out'] = (out,) + native_method = getattr(native_ufunc, method, None) + use_native = (use_native and native_method is not None) + + if use_native: + # Native method could exist but raise `NotImplementedError` + # or return `NotImplemented`, falling back to base case + # then, too + try: + res = native_method(*inputs, **kwargs) + except NotImplementedError: + res = super(CupyTensor, self).__array_ufunc__( + ufunc, method, *inputs, **kwargs) + else: + if res is NotImplemented: + res = super(CupyTensor, self).__array_ufunc__( + ufunc, method, *inputs, **kwargs) + + else: + res = super(CupyTensor, self).__array_ufunc__( + ufunc, method, *inputs, **kwargs) + + # Shortcut for scalar or no return value + if np.isscalar(res) or res is None: + # The first occurs for `reduce` with all axes, + # the second for in-place stuff (`at` currently) + return res + + # Wrap result if necessary (lazily) + if out is None: + if is_floating_dtype(res.dtype): + if res.shape != self.shape: + # Don't propagate weighting if shape changes + weighting = CupyTensorSpaceConstWeighting(1.0, + exponent) + spc_kwargs = {'weighting': weighting} + else: + spc_kwargs = {} + + res_space = type(self.space)( + res.shape, res.dtype, self.device, **spc_kwargs) + result = res_space.element(res) + else: + result = out_tuple[0] + + return result + + @property + def ufuncs(self): + """Access to NumPy style ufuncs. + + Examples + -------- + >>> r2 = odl.rn(2, impl='cupy') + >>> x = r2.element([1, -2]) + >>> x.ufuncs.absolute() + rn(2, impl='cupy').element([ 1., 2.]) + + These functions can also be used with broadcasting or + array-like input: + + >>> x.ufuncs.add(3) + rn(2, impl='cupy').element([ 4., 1.]) + >>> x.ufuncs.subtract([3, 3]) + rn(2, impl='cupy').element([-2., -5.]) + + There is also support for various reductions + (sum, prod, amin, amax): + + >>> x.ufuncs.sum() + -1.0 + >>> x.ufuncs.prod() + -2.0 + + They also support an out parameter + + >>> y = r2.element([3, 4]) + >>> out = r2.element() + >>> result = x.ufuncs.add(y, out=out) + >>> result + rn(2, impl='cupy').element([ 4., 2.]) + >>> result is out + True + + Notes + ----- + Those ufuncs which are implemented natively on the GPU incur no + significant overhead. However, for missing functions, a fallback + Numpy implementation is used which causes significant overhead + due to data copies between host and device. + """ + # TODO: Test with some native ufuncs, then remove this attribute + return super(CupyTensor, self).ufuncs + + @property + def real(self): + """Real part of this tensor. + + Returns + ------- + real : `CupyTensor` view with real dtype + The real part of this tensor as an element of an `rn` space. + """ + # Only real dtypes currently + return self + + @real.setter + def real(self, newreal): + """Setter for the real part. + + This method is invoked by ``tensor.real = other``. + + Parameters + ---------- + newreal : `array-like` or scalar + The new real part for this tensor. + """ + self.real.data[:] = newreal + + @property + def imag(self): + """Imaginary part of this tensor. + + Returns + ------- + imag : `CupyTensor` + The imaginary part of this tensor as an element of an `rn` space. + """ + # Only real dtypes currently + return self.space.zero() + + @imag.setter + def imag(self, newimag): + """Setter for the imaginary part. + + This method is invoked by ``tensor.imag = other``. + + Parameters + ---------- + newimag : `array-like` or scalar + The new imaginary part for this tensor. + """ + raise NotImplementedError('complex dtypes not supported') + + def conj(self, out=None): + """Complex conjugate of this tensor. + + Parameters + ---------- + out : `CupyTensor`, optional + Tensor to which the complex conjugate is written. + Must be an element of this tensor's space. + + Returns + ------- + out : `CupyTensor` + The complex conjugate tensor. If ``out`` was provided, + the returned object is a reference to it. + """ + # Only real dtypes currently + if out is None: + return self.copy() + else: + self.assign(out) + return out + + def __ipow__(self, other): + """Return ``self **= other``.""" + try: + if other == int(other): + return super(CupyTensorSpace, self).__ipow__(other) + except TypeError: + pass + + self.ufuncs.power(self.data, other, out=self.data) + return self + + +# --- Weightings --- # + + +def _weighting(weights, exponent): + """Return a weighting whose type is inferred from the arguments.""" + if np.isscalar(weights): + weighting = CupyTensorSpaceConstWeighting(weights, exponent=exponent) + else: + # TODO: sequence of 1D array-likes + weights = cupy.array(weights, copy=False) + weighting = CupyTensorSpaceArrayWeighting(weights, exponent=exponent) + return weighting + + +# Kernels for space functions + +dotw = cupy.ReductionKernel(in_params='T x, T y, W w', + out_params='T res', + map_expr='x * y * w', + reduce_expr='a + b', + post_map_expr='res = a', + identity='0', + name='dotw') + +nrm0 = cupy.ReductionKernel(in_params='T x', + out_params='int64 res', + map_expr='x != 0', + reduce_expr='a + b', + post_map_expr='res = a', + identity='0', + name='nrm0') + +nrm1 = cupy.ReductionKernel(in_params='T x', + out_params='T res', + map_expr='abs(x)', + reduce_expr='a + b', + post_map_expr='res = a', + identity='0', + name='nrm1w') + +nrm1w = cupy.ReductionKernel(in_params='T x, W w', + out_params='T res', + map_expr='abs(x) * w', + reduce_expr='a + b', + post_map_expr='res = a', + identity='0', + name='nrm1w') + +nrm2 = cupy.ReductionKernel(in_params='T x', + out_params='T res', + map_expr='x * x', + reduce_expr='a + b', + post_map_expr='res = sqrt(a)', + identity='0', + name='nrm2') + +nrm2w = cupy.ReductionKernel(in_params='T x, W w', + out_params='T res', + map_expr='x * x * w', + reduce_expr='a + b', + post_map_expr='res = sqrt(a)', + identity='0', + name='nrm2w') + +nrminf = cupy.ReductionKernel(in_params='T x', + out_params='T res', + map_expr='abs(x)', + reduce_expr='a > b ? a : b', + post_map_expr='res = a', + identity='0', + name='nrminf') + +nrmneginf = cupy.ReductionKernel(in_params='T x', + out_params='T res', + map_expr='abs(x)', + reduce_expr='a > b ? b : a', + post_map_expr='res = a', + identity='0', + name='nrmneginf') + +nrmp = cupy.ReductionKernel(in_params='T x, T p', + out_params='T res', + map_expr='pow(abs(x), p)', + reduce_expr='a + b', + post_map_expr='res = pow(a, 1 / p)', + identity='0', + name='nrmp') + +nrmpw = cupy.ReductionKernel(in_params='T x, T p, W w', + out_params='T res', + map_expr='pow(abs(x), p) * w', + reduce_expr='a + b', + post_map_expr='res = pow(a, 1 / p)', + identity='0', + name='nrmpw') + +dist0 = cupy.ReductionKernel(in_params='T x, T y', + out_params='int64 res', + map_expr='x != y', + reduce_expr='a + b', + post_map_expr='res = a', + identity='0', + name='dist0') + +dist1 = cupy.ReductionKernel(in_params='T x, T y', + out_params='T res', + map_expr='abs(x - y)', + reduce_expr='a + b', + post_map_expr='res = a', + identity='0', + name='dist1') + +dist1w = cupy.ReductionKernel(in_params='T x, T y, W w', + out_params='T res', + map_expr='abs(x - y) * w', + reduce_expr='a + b', + post_map_expr='res = a', + identity='0', + name='dist1w') + +dist2 = cupy.ReductionKernel(in_params='T x, T y', + out_params='T res', + map_expr='(x - y) * (x - y)', + reduce_expr='a + b', + post_map_expr='res = sqrt(a)', + identity='0', + name='dist2') + +dist2w = cupy.ReductionKernel(in_params='T x, T y, W w', + out_params='T res', + map_expr='(x - y) * (x - y) * w', + reduce_expr='a + b', + post_map_expr='res = sqrt(a)', + identity='0', + name='dist2w') + +distinf = cupy.ReductionKernel(in_params='T x, T y', + out_params='T res', + map_expr='abs(x - y)', + reduce_expr='a > b ? a : b', + post_map_expr='res = a', + identity='0', + name='distinf') + +distneginf = cupy.ReductionKernel(in_params='T x, T y', + out_params='T res', + map_expr='abs(x - y)', + reduce_expr='a > b ? b : a', + post_map_expr='res = a', + identity='0', + name='distneginf') + +distp = cupy.ReductionKernel(in_params='T x, T y, T p', + out_params='T res', + map_expr='pow(abs(x - y), p)', + reduce_expr='a + b', + post_map_expr='res = pow(a, 1 / p)', + identity='0', + name='distp') + +distpw = cupy.ReductionKernel(in_params='T x, T y, T p, W w', + out_params='T res', + map_expr='pow(abs(x - y), p) * w', + reduce_expr='a + b', + post_map_expr='res = pow(a, 1 / p)', + identity='0', + name='distpw') + + +class CupyTensorSpaceArrayWeighting(ArrayWeighting): + + """Array weighting for `CupyTensorSpace`. + + See `ArrayWeighting` for further details. + """ + + def __init__(self, array, exponent=2.0): + """Initialize a new instance. + + Parameters + ---------- + array : `array-like`, one-dim. + Weighting array of the inner product, norm and distance. + exponent : positive float + Exponent of the norm. For values other than 2.0, the inner + product is not defined. + """ + if isinstance(array, CupyTensor): + array = array.data + elif not isinstance(array, cupy.ndarray): + array = cupy.array(array, copy=False) + super(CupyTensorSpaceArrayWeighting, self).__init__( + array, impl='cupy', exponent=exponent) + + def inner(self, x1, x2): + """Return the weighted inner product of two tensors. + + Parameters + ---------- + x1, x2 : `CupyTensor` + Tensors whose inner product is calculated. + + Returns + ------- + inner : float or complex + The inner product of the two provided tensors. + """ + if self.exponent != 2.0: + raise NotImplementedError('no inner product defined for ' + 'exponent != 2 (got {})' + ''.format(self.exponent)) + else: + return x1.space.field.element(dotw(x1.data, x2.data, self.array)) + + def norm(self, x): + """Return the weighted norm of a tensor. + + Parameters + ---------- + x : `CupyTensor` + Tensor whose norm is calculated. + + Returns + ------- + norm : float + The norm of the provided tensor. + """ + if self.exponent == 0: + return float(nrm0(x.data)) + elif self.exponent == 1: + return float(nrm1w(x.data, self.array)) + elif self.exponent == 2: + return float(nrm2w(x.data, self.array)) + elif self.exponent == float('inf'): + return float(nrminf(x.data)) + elif self.exponent == -float('inf'): + return float(nrmneginf(x.data)) + else: + return float(nrmpw(x.data, self.exponent, self.array)) + + def dist(self, x1, x2): + """Return the weighted distance of two tensors. + + Parameters + ---------- + x1, x2 : `CupyTensor` + Tensors whose mutual distance is calculated. + + Returns + ------- + dist : float + The distance between the provided tensors. + """ + if self.exponent == 0: + return float(dist0(x1.data, x2.data)) + elif self.exponent == 1: + return float(dist1w(x1.data, x2.data, self.array)) + elif self.exponent == 2: + return float(dist2w(x1.data, x2.data, self.array)) + elif self.exponent == float('inf'): + return float(distinf(x1.data, x2.data, self.array)) + elif self.exponent == -float('inf'): + return float(distneginf(x1.data, x2.data, self.array)) + else: + return float(distpw(x1.data, x2.data, self.exponent, self.array)) + + +class CupyTensorSpaceConstWeighting(ConstWeighting): + + """Constant weighting for `CupyTensorSpace`. + + See `ConstWeighting` for further details. + """ + + def __init__(self, constant, exponent=2.0): + """Initialize a new instance. + + Parameters + ---------- + constant : positive float + Weighting constant of the inner product. + exponent : positive float + Exponent of the norm. For values other than 2.0, the inner + product is not defined. + """ + super(CupyTensorSpaceConstWeighting, self).__init__( + constant, impl='cupy', exponent=exponent) + + def inner(self, x1, x2): + """Return the weighted inner product of two tensors. + + Parameters + ---------- + x1, x2 : `CupyTensor` + Tensors whose inner product is calculated. + + Returns + ------- + inner : float or complex + The inner product of the two provided tensors. + """ + if self.exponent != 2.0: + raise NotImplementedError('no inner product defined for ' + 'exponent != 2 (got {})' + ''.format(self.exponent)) + else: + return x1.space.field.element( + self.const * cupy.dot(x1.data, x2.data)) + + def norm(self, x): + """Return the constant-weighted norm of a tensor. + + Parameters + ---------- + x1 : `CupyTensor` + Tensor whose norm is calculated. + + Returns + ------- + norm : float + The norm of the tensor. + """ + if self.exponent == 0: + return float(nrm0(x.data)) + elif self.exponent == 1: + return float(self.const * nrm1(x.data)) + elif self.exponent == 2: + # We try to use cuBLAS nrm2 + try: + incx = _flat_inc(x.data) + except ValueError: + use_cublas = False + else: + use_cublas = True + + if use_cublas: + try: + nrm2_cublas = _cublas_func('nrm2', x.dtype) + except (ValueError, AttributeError): + pass + else: + with cupy.cuda.Device(x.device) as dev: + norm = nrm2_cublas( + dev.cublas_handle, x.size, x.data_ptr, incx) + return float(np.sqrt(self.const) * norm) + + # Cannot use cuBLAS, fall back to custom kernel + return float(np.sqrt(self.const) * nrm2(x.data)) + elif self.exponent == float('inf'): + return float(nrminf(x.data)) + elif self.exponent == -float('inf'): + return float(nrmneginf(x.data)) + else: + return float(self.const ** (1 / self.exponent) * + nrmp(x, self.exponent)) + + def dist(self, x1, x2): + """Return the weighted distance between two tensors. + + Parameters + ---------- + x1, x2 : `CupyTensor` + Tensors whose mutual distance is calculated. + + Returns + ------- + dist : float + The distance between the tensors. + """ + if self.exponent == 0: + return float(dist0(x1.data, x2.data)) + elif self.exponent == 1: + return float(self.const * dist1(x1.data, x2.data)) + elif self.exponent == 2: + # cuBLAS nrm2(x1 - x2) would probably be faster, but would + # require a copy, so we don't do it + return float(np.sqrt(self.const) * dist2(x1.data, x2.data)) + elif self.exponent == float('inf'): + return float(distinf(x1.data, x2.data)) + elif self.exponent == -float('inf'): + return float(distneginf(x1.data, x2.data)) + else: + return float(self.const ** (1 / self.exponent) * + distp(x1.data, x2.data, self.exponent)) + + +class CupyTensorSpaceCustomInner(CustomInner): + + """Class for handling custom inner products in `CupyTensorSpace`.""" + + def __init__(self, inner): + """Initialize a new instance. + + Parameters + ---------- + inner : callable + The inner product implementation. It must accept two + `CupyTensor` arguments, return an element from their space's + field (real or complex number) and satisfy the following + conditions for all vectors ``x, y, z`` and scalars ``s``: + + - `` = conj()`` + - `` = s * + `` + - `` = 0`` if and only if ``x = 0`` + """ + super(CupyTensorSpaceCustomInner, self).__init__(inner, impl='cupy') + + +class CupyTensorSpaceCustomNorm(CustomNorm): + + """Class for handling a user-specified norm in `CupyTensorSpace`. + + Note that this removes ``inner``. + """ + + def __init__(self, norm): + """Initialize a new instance. + + Parameters + ---------- + norm : callable + The norm implementation. It must accept an `CupyTensor` + argument, return a float and satisfy the following + conditions for all vectors ``x, y`` and scalars ``s``: + + - ``||x|| >= 0`` + - ``||x|| = 0`` if and only if ``x = 0`` + - ``||s * x|| = |s| * ||x||`` + - ``||x + y|| <= ||x|| + ||y||`` + """ + super(CupyTensorSpaceCustomNorm, self).__init__(norm, impl='cupy') + + +class CupyTensorSpaceCustomDist(CustomDist): + + """Class for handling a user-specified distance in `CupyTensorSpace`. + + Note that this removes ``inner`` and ``norm``. + """ + + def __init__(self, dist): + """Initialize a new instance. + + Parameters + ---------- + dist : callable + The distance function defining a metric on `CupyTensorSpace`. + It must accept two `CupyTensor` arguments, return a float and + fulfill the following mathematical conditions for any three + vectors ``x, y, z``: + + - ``dist(x, y) >= 0`` + - ``dist(x, y) = 0`` if and only if ``x = y`` + - ``dist(x, y) = dist(y, x)`` + - ``dist(x, y) <= dist(x, z) + dist(z, y)`` + """ + super(CupyTensorSpaceCustomDist, self).__init__(dist, impl='cupy') + + +if __name__ == '__main__': + if CUPY_AVAILABLE: + from odl.util.testutils import run_doctests + run_doctests() diff --git a/odl/space/entry_points.py b/odl/space/entry_points.py index 1c610d4b9f5..af28ed9a97a 100644 --- a/odl/space/entry_points.py +++ b/odl/space/entry_points.py @@ -23,12 +23,15 @@ from __future__ import print_function, division, absolute_import from odl.space.npy_tensors import NumpyTensorSpace +from odl.space.cupy_tensors import CUPY_AVAILABLE, CupyTensorSpace # We don't expose anything to odl.space __all__ = () IS_INITIALIZED = False TENSOR_SPACE_IMPLS = {'numpy': NumpyTensorSpace} +if CUPY_AVAILABLE: + TENSOR_SPACE_IMPLS['cupy'] = CupyTensorSpace def _initialize_if_needed(): From f8fed7afc7c6c3ae503ae3b51d37e0830e6c40e9 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Thu, 23 Nov 2017 11:56:18 +0100 Subject: [PATCH 02/38] MAINT: improve import of cupy --- odl/space/cupy_tensors.py | 330 ++++++++++++++++++++------------------ 1 file changed, 173 insertions(+), 157 deletions(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index 27480cd4543..0a89eae2b74 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -6,10 +6,15 @@ # v. 2.0. If a copy of the MPL was not distributed with this file, You can # obtain one at https://mozilla.org/MPL/2.0/. -"""Implementation of tensor spaces using ``pygpu``.""" +"""Implementation of tensor spaces using CuPy. + +See https://cupy.chainer.org/ or https://github.com/cupy/cupy for details +on the backend. +""" from __future__ import print_function, division, absolute_import import numpy as np +import warnings from odl.set import RealNumbers from odl.space.base_tensors import TensorSpace, Tensor @@ -23,9 +28,11 @@ except ImportError: CUPY_AVAILABLE = False else: - from pkg_resources import parse_version - if parse_version(cupy.__version__) < parse_version('2.0.0rc1'): - raise ImportError('cupy <2.0.0rc1 not supported') + _maj = int(cupy.__version__.split('.')[0]) + if _maj < 2: + raise warnings.warn( + 'your version {} of CuPy is not supported; please upgrade to ' + 'version 2.0.0 or higher'.format(cupy.__version__), RuntimeWarning) CUPY_AVAILABLE = True @@ -35,10 +42,11 @@ # --- Space method implementations --- # -lico = cupy.ElementwiseKernel(in_params='T a, T x, T b, T y', - out_params='T z', - operation='z = a * x + b * y;', - name='lico') +if CUPY_AVAILABLE: + lico = cupy.ElementwiseKernel(in_params='T a, T x, T b, T y', + out_params='T z', + operation='z = a * x + b * y;', + name='lico') def fallback_scal(a, x): @@ -77,12 +85,17 @@ def _cublas_func(name, dtype): ValueError : If the data type is not supported by cuBLAS. """ - if np.dtype(dtype) == 'float32': + dtype, dtype_in = np.dtype(dtype), dtype + if dtype == 'float32': prefix = 's' - elif np.dtype(dtype) == 'float64': + elif dtype == 'float64': prefix = 'd' + elif dtype == 'complex64': + prefix = 'c' + elif dtype == 'complex128': + prefix = 'z' else: - raise ValueError('dtype {!r} not supported by cuBLAS'.format(dtype)) + raise ValueError('dtype {!r} not supported by cuBLAS'.format(dtype_in)) return getattr(cupy.cuda.cublas, prefix + name) @@ -197,13 +210,15 @@ def _lincomb_impl(a, x1, b, x2, out): class CupyTensorSpace(TensorSpace): - """Tensor space implemented with GPU arrays. + """Tensor space implemented with CUDA arrays using the CuPy library. This space implements tensors of arbitrary rank over a `Field` ``F``, which is either the real or complex numbers. Its elements are represented as instances of the `CupyTensor` class. + + See https://github.com/cupy/cupy for details on the backend. """ def __init__(self, shape, dtype='float64', device=None, **kwargs): @@ -1799,157 +1814,158 @@ def _weighting(weights, exponent): # Kernels for space functions -dotw = cupy.ReductionKernel(in_params='T x, T y, W w', - out_params='T res', - map_expr='x * y * w', - reduce_expr='a + b', - post_map_expr='res = a', - identity='0', - name='dotw') - -nrm0 = cupy.ReductionKernel(in_params='T x', - out_params='int64 res', - map_expr='x != 0', - reduce_expr='a + b', - post_map_expr='res = a', - identity='0', - name='nrm0') - -nrm1 = cupy.ReductionKernel(in_params='T x', - out_params='T res', - map_expr='abs(x)', - reduce_expr='a + b', - post_map_expr='res = a', - identity='0', - name='nrm1w') - -nrm1w = cupy.ReductionKernel(in_params='T x, W w', - out_params='T res', - map_expr='abs(x) * w', - reduce_expr='a + b', - post_map_expr='res = a', - identity='0', - name='nrm1w') - -nrm2 = cupy.ReductionKernel(in_params='T x', - out_params='T res', - map_expr='x * x', - reduce_expr='a + b', - post_map_expr='res = sqrt(a)', - identity='0', - name='nrm2') - -nrm2w = cupy.ReductionKernel(in_params='T x, W w', - out_params='T res', - map_expr='x * x * w', - reduce_expr='a + b', - post_map_expr='res = sqrt(a)', - identity='0', - name='nrm2w') - -nrminf = cupy.ReductionKernel(in_params='T x', - out_params='T res', - map_expr='abs(x)', - reduce_expr='a > b ? a : b', - post_map_expr='res = a', - identity='0', - name='nrminf') - -nrmneginf = cupy.ReductionKernel(in_params='T x', +if CUPY_AVAILABLE: + dotw = cupy.ReductionKernel(in_params='T x, T y, W w', + out_params='T res', + map_expr='x * y * w', + reduce_expr='a + b', + post_map_expr='res = a', + identity='0', + name='dotw') + + nrm0 = cupy.ReductionKernel(in_params='T x', + out_params='int64 res', + map_expr='x != 0', + reduce_expr='a + b', + post_map_expr='res = a', + identity='0', + name='nrm0') + + nrm1 = cupy.ReductionKernel(in_params='T x', + out_params='T res', + map_expr='abs(x)', + reduce_expr='a + b', + post_map_expr='res = a', + identity='0', + name='nrm1w') + + nrm1w = cupy.ReductionKernel(in_params='T x, W w', out_params='T res', - map_expr='abs(x)', - reduce_expr='a > b ? b : a', + map_expr='abs(x) * w', + reduce_expr='a + b', post_map_expr='res = a', identity='0', - name='nrmneginf') - -nrmp = cupy.ReductionKernel(in_params='T x, T p', - out_params='T res', - map_expr='pow(abs(x), p)', - reduce_expr='a + b', - post_map_expr='res = pow(a, 1 / p)', - identity='0', - name='nrmp') - -nrmpw = cupy.ReductionKernel(in_params='T x, T p, W w', - out_params='T res', - map_expr='pow(abs(x), p) * w', - reduce_expr='a + b', - post_map_expr='res = pow(a, 1 / p)', - identity='0', - name='nrmpw') - -dist0 = cupy.ReductionKernel(in_params='T x, T y', - out_params='int64 res', - map_expr='x != y', - reduce_expr='a + b', - post_map_expr='res = a', - identity='0', - name='dist0') - -dist1 = cupy.ReductionKernel(in_params='T x, T y', - out_params='T res', - map_expr='abs(x - y)', - reduce_expr='a + b', - post_map_expr='res = a', - identity='0', - name='dist1') - -dist1w = cupy.ReductionKernel(in_params='T x, T y, W w', - out_params='T res', - map_expr='abs(x - y) * w', - reduce_expr='a + b', - post_map_expr='res = a', - identity='0', - name='dist1w') - -dist2 = cupy.ReductionKernel(in_params='T x, T y', - out_params='T res', - map_expr='(x - y) * (x - y)', - reduce_expr='a + b', - post_map_expr='res = sqrt(a)', - identity='0', - name='dist2') - -dist2w = cupy.ReductionKernel(in_params='T x, T y, W w', - out_params='T res', - map_expr='(x - y) * (x - y) * w', - reduce_expr='a + b', - post_map_expr='res = sqrt(a)', - identity='0', - name='dist2w') - -distinf = cupy.ReductionKernel(in_params='T x, T y', - out_params='T res', - map_expr='abs(x - y)', - reduce_expr='a > b ? a : b', - post_map_expr='res = a', - identity='0', - name='distinf') - -distneginf = cupy.ReductionKernel(in_params='T x, T y', + name='nrm1w') + + nrm2 = cupy.ReductionKernel(in_params='T x', + out_params='T res', + map_expr='x * x', + reduce_expr='a + b', + post_map_expr='res = sqrt(a)', + identity='0', + name='nrm2') + + nrm2w = cupy.ReductionKernel(in_params='T x, W w', + out_params='T res', + map_expr='x * x * w', + reduce_expr='a + b', + post_map_expr='res = sqrt(a)', + identity='0', + name='nrm2w') + + nrminf = cupy.ReductionKernel(in_params='T x', out_params='T res', - map_expr='abs(x - y)', - reduce_expr='a > b ? b : a', + map_expr='abs(x)', + reduce_expr='a > b ? a : b', post_map_expr='res = a', identity='0', - name='distneginf') - -distp = cupy.ReductionKernel(in_params='T x, T y, T p', - out_params='T res', - map_expr='pow(abs(x - y), p)', - reduce_expr='a + b', - post_map_expr='res = pow(a, 1 / p)', - identity='0', - name='distp') - -distpw = cupy.ReductionKernel(in_params='T x, T y, T p, W w', - out_params='T res', - map_expr='pow(abs(x - y), p) * w', - reduce_expr='a + b', - post_map_expr='res = pow(a, 1 / p)', - identity='0', - name='distpw') + name='nrminf') + + nrmneginf = cupy.ReductionKernel(in_params='T x', + out_params='T res', + map_expr='abs(x)', + reduce_expr='a > b ? b : a', + post_map_expr='res = a', + identity='0', + name='nrmneginf') + + nrmp = cupy.ReductionKernel(in_params='T x, T p', + out_params='T res', + map_expr='pow(abs(x), p)', + reduce_expr='a + b', + post_map_expr='res = pow(a, 1 / p)', + identity='0', + name='nrmp') + + nrmpw = cupy.ReductionKernel(in_params='T x, T p, W w', + out_params='T res', + map_expr='pow(abs(x), p) * w', + reduce_expr='a + b', + post_map_expr='res = pow(a, 1 / p)', + identity='0', + name='nrmpw') + + dist0 = cupy.ReductionKernel(in_params='T x, T y', + out_params='int64 res', + map_expr='x != y', + reduce_expr='a + b', + post_map_expr='res = a', + identity='0', + name='dist0') + + dist1 = cupy.ReductionKernel(in_params='T x, T y', + out_params='T res', + map_expr='abs(x - y)', + reduce_expr='a + b', + post_map_expr='res = a', + identity='0', + name='dist1') + + dist1w = cupy.ReductionKernel(in_params='T x, T y, W w', + out_params='T res', + map_expr='abs(x - y) * w', + reduce_expr='a + b', + post_map_expr='res = a', + identity='0', + name='dist1w') + + dist2 = cupy.ReductionKernel(in_params='T x, T y', + out_params='T res', + map_expr='(x - y) * (x - y)', + reduce_expr='a + b', + post_map_expr='res = sqrt(a)', + identity='0', + name='dist2') + + dist2w = cupy.ReductionKernel(in_params='T x, T y, W w', + out_params='T res', + map_expr='(x - y) * (x - y) * w', + reduce_expr='a + b', + post_map_expr='res = sqrt(a)', + identity='0', + name='dist2w') + + distinf = cupy.ReductionKernel(in_params='T x, T y', + out_params='T res', + map_expr='abs(x - y)', + reduce_expr='a > b ? a : b', + post_map_expr='res = a', + identity='0', + name='distinf') + + distneginf = cupy.ReductionKernel(in_params='T x, T y', + out_params='T res', + map_expr='abs(x - y)', + reduce_expr='a > b ? b : a', + post_map_expr='res = a', + identity='0', + name='distneginf') + + distp = cupy.ReductionKernel(in_params='T x, T y, T p', + out_params='T res', + map_expr='pow(abs(x - y), p)', + reduce_expr='a + b', + post_map_expr='res = pow(a, 1 / p)', + identity='0', + name='distp') + + distpw = cupy.ReductionKernel(in_params='T x, T y, T p, W w', + out_params='T res', + map_expr='pow(abs(x - y), p) * w', + reduce_expr='a + b', + post_map_expr='res = pow(a, 1 / p)', + identity='0', + name='distpw') class CupyTensorSpaceArrayWeighting(ArrayWeighting): From 1c64fcb3dbba80de6cb1e9c991b88625871537e8 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Thu, 23 Nov 2017 18:59:30 +0100 Subject: [PATCH 03/38] WIP: fix cupy->numpy transfers --- odl/space/cupy_tensors.py | 140 +++++++++++++++++++++++++++----------- 1 file changed, 101 insertions(+), 39 deletions(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index 0a89eae2b74..9577daecd50 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -16,12 +16,13 @@ import numpy as np import warnings -from odl.set import RealNumbers +from odl.set import RealNumbers, ComplexNumbers from odl.space.base_tensors import TensorSpace, Tensor from odl.space.weighting import ( Weighting, ArrayWeighting, ConstWeighting, CustomInner, CustomNorm, CustomDist) -from odl.util import dtype_str, is_floating_dtype, signature_string +from odl.util import ( + array_str, dtype_str, is_floating_dtype, signature_string, indent) try: import cupy @@ -599,13 +600,14 @@ def __eq__(self, other): >>> same_space == space True """ - return (super().__eq__(other) and + return (super(CupyTensorSpace, self).__eq__(other) and self.device == other.device and self.weighting == other.weighting) def __hash__(self): """Return ``hash(self)``.""" - return hash((super().__hash__(), self.device, self.weighting)) + return hash((super(CupyTensorSpace, self).__hash__(), self.device, + self.weighting)) def _lincomb(self, a, x1, b, x2, out): """Linear combination of ``x1`` and ``x2``. @@ -874,12 +876,16 @@ def default_dtype(field=None): dtype : `numpy.dtype` Numpy data type specifier. The returned defaults are: - ``RealNumbers()`` : ``np.dtype('float64')`` + - ``RealNumbers()`` or ``None`` : ``np.dtype('float64')`` + - ``ComplexNumbers()`` : ``np.dtype('complex128')`` - ``ComplexNumbers()`` : not supported + These choices correspond to the defaults of the ``cupy`` + library. """ if field is None or field == RealNumbers(): return np.dtype('float64') + elif field == ComplexNumbers(): + return np.dtype('complex128') else: raise ValueError('no default data type defined for field {}.' ''.format(field)) @@ -1029,7 +1035,7 @@ def copy(self): Returns ------- - copy : `pygpu._array.ndgpuarray` + copy : `CupyTensor` A deep copy. Examples @@ -1056,7 +1062,7 @@ def __getitem__(self, indices): Returns ------- - values : scalar or `pygpu._array.ndgpuarray` + values : scalar or `cupy.core.core.ndarray` The value(s) at the index (indices). Examples @@ -1107,11 +1113,11 @@ def __getitem__(self, indices): arr = self.data[indices] if arr.shape == (): if arr.dtype.kind == 'f': - return float(np.asarray(arr)) + return float(cupy.asnumpy(arr)) elif arr.dtype.kind == 'c': - return complex(np.asarray(arr)) + return complex(cupy.asnumpy(arr)) elif arr.dtype.kind in ('u', 'i'): - return int(np.asarray(arr)) + return int(cupy.asnumpy(arr)) else: raise RuntimeError("no conversion for dtype {}" "".format(arr.dtype)) @@ -1280,12 +1286,22 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): for further details. See also the `general documentation on Numpy ufuncs`_. - .. note:: - This implementation looks for native ufuncs in ``pygpu.ufuncs`` - and falls back to the basic implementation with Numpy arrays - in case no native ufunc is available. That fallback version - comes with significant overhead due to data copies between - host and device. + .. warning:: + Apart from ``'__call__'`` (invoked by, e.g., ``np.add(x, y))``, + CuPy has no native implementation of ufunc methods like + ``'reduce'`` or ``'accumulate'``. We manually implement the + mappings (covering most use cases) + + - ``np.add.reduce`` -> ``cupy.sum`` + - ``np.add.accumulate`` -> ``cupy.cumsum`` + - ``np.multiply.reduce`` -> ``cupy.prod`` + - ``np.multiply.reduce`` -> ``cupy.cumprod``. + + **All other such methods will run Numpy code and be slow**! + + Please consult the `CuPy documentation on ufuncs + `_ + to check the current state of the library. .. note:: When an ``out`` parameter is specified, and (one of) it has @@ -1507,10 +1523,10 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): inp.data if isinstance(inp, type(self)) else inp for inp in inputs) - # For native ufuncs, we turn non-scalar inputs into cupy arrays, - # as a workaround for https://github.com/cupy/cupy/issues/594 - # TODO: remove code when the upstream issue is fixed if use_native: + # TODO: remove when upstream issue is fixed + # For native ufuncs, we turn non-scalar inputs into cupy arrays, + # as a workaround for https://github.com/cupy/cupy/issues/594 inputs, orig_inputs = [], inputs for inp in orig_inputs: if (isinstance(inp, cupy.ndarray) or @@ -1519,6 +1535,20 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): inputs.append(inp) else: inputs.append(cupy.array(inp)) + elif method != 'at': + # TODO: remove when upstream issue is fixed + # For non-native ufuncs (except `at`), we need ot cast our tensors + # and Cupy arrays to Numpy arrays explicitly, since `__array__` + # and friends are not implemented. See + # https://github.com/cupy/cupy/issues/589 + inputs, orig_inputs = [], inputs + for inp in orig_inputs: + if isinstance(inp, cupy.ndarray): + inputs.append(cupy.asnumpy(inp)) + elif isinstance(inp, CupyTensor): + inputs.append(cupy.asnumpy(inp.data)) + else: + inputs.append(inp) # --- Get some parameters for later --- # @@ -1595,20 +1625,19 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): def eval_at_via_npy(*inputs, **kwargs): import ctypes cupy_arr = inputs[0] - npy_arr = np.asarray(cupy_arr) + npy_arr = cupy.asnumpy(cupy_arr) new_inputs = (npy_arr,) + inputs[1:] super(CupyTensor, self).__array_ufunc__( ufunc, method, *new_inputs, **kwargs) # Workaround for https://github.com/cupy/cupy/issues/593 - # TODO: use cupy_arr[:] = npy_arr when it's fixed and not - # slower + # TODO: use cupy_arr[:] = npy_arr when available cupy_arr.data.copy_from_host( npy_arr.ctypes.data_as(ctypes.c_void_p), npy_arr.nbytes) if use_native: # Native method could exist but raise `NotImplementedError` - # or return `NotImplemented`, falling back to Numpy case - # then, too + # or return `NotImplemented`. We fall back to Numpy also in + # that situation. try: res = native_method(*inputs, **kwargs) except NotImplementedError: @@ -1626,8 +1655,8 @@ def eval_at_via_npy(*inputs, **kwargs): if use_native: # Native method could exist but raise `NotImplementedError` - # or return `NotImplemented`, falling back to base case - # then, too + # or return `NotImplemented`. We fall back to Numpy also in + # that situation. try: res = native_method(*inputs, **kwargs) except NotImplementedError: @@ -1653,8 +1682,8 @@ def eval_at_via_npy(*inputs, **kwargs): if is_floating_dtype(res.dtype): if res.shape != self.shape: # Don't propagate weighting if shape changes - weighting = CupyTensorSpaceConstWeighting(1.0, - exponent) + weighting = CupyTensorSpaceConstWeighting( + 1.0, exponent) spc_kwargs = {'weighting': weighting} else: spc_kwargs = {} @@ -1723,8 +1752,10 @@ def real(self): real : `CupyTensor` view with real dtype The real part of this tensor as an element of an `rn` space. """ - # Only real dtypes currently - return self + if self.space.is_real: + return self + else: + return self.space.real_space.element(self.data.real) @real.setter def real(self, newreal): @@ -1737,7 +1768,7 @@ def real(self, newreal): newreal : `array-like` or scalar The new real part for this tensor. """ - self.real.data[:] = newreal + self.data.real[:] = newreal @property def imag(self): @@ -1748,8 +1779,10 @@ def imag(self): imag : `CupyTensor` The imaginary part of this tensor as an element of an `rn` space. """ - # Only real dtypes currently - return self.space.zero() + if self.space.is_real: + return self.space.zero() + else: + return self.space.real_space.element(self.data.imag) @imag.setter def imag(self, newimag): @@ -1762,7 +1795,7 @@ def imag(self, newimag): newimag : `array-like` or scalar The new imaginary part for this tensor. """ - raise NotImplementedError('complex dtypes not supported') + self.data.imag[:] = newimag def conj(self, out=None): """Complex conjugate of this tensor. @@ -1779,11 +1812,17 @@ def conj(self, out=None): The complex conjugate tensor. If ``out`` was provided, the returned object is a reference to it. """ - # Only real dtypes currently if out is None: - return self.copy() + if self.space.is_real: + return self.copy() + else: + return self.space.element(self.data.conj()) else: - self.assign(out) + if self.space.is_real: + self.assign(out) + else: + # In-place not available as it seems + out[:] = self.data.conj() return out def __ipow__(self, other): @@ -1806,7 +1845,8 @@ def _weighting(weights, exponent): if np.isscalar(weights): weighting = CupyTensorSpaceConstWeighting(weights, exponent=exponent) else: - # TODO: sequence of 1D array-likes + # TODO: sequence of 1D array-likes, see + # https://github.com/odlgroup/odl/pull/1238 weights = cupy.array(weights, copy=False) weighting = CupyTensorSpaceArrayWeighting(weights, exponent=exponent) return weighting @@ -2065,6 +2105,28 @@ def dist(self, x1, x2): else: return float(distpw(x1.data, x2.data, self.exponent, self.array)) + # TODO: remove repr_part and __repr__ when cupy.ndarray.__array__ + # is implemented. See + # https://github.com/cupy/cupy/issues/589 + @property + def repr_part(self): + """String usable in a space's ``__repr__`` method.""" + # TODO: use edgeitems + arr_str = array_str(cupy.asnumpy(self.array), nprint=10) + optargs = [('weighting', arr_str, ''), + ('exponent', self.exponent, 2.0)] + return signature_string([], optargs, sep=',\n', + mod=[[], ['!s', ':.4']]) + + def __repr__(self): + """Return ``repr(self)``.""" + # TODO: use edgeitems + posargs = [array_str(cupy.asnumpy(self.array), nprint=10)] + optargs = [('exponent', self.exponent, 2.0)] + inner_str = signature_string(posargs, optargs, sep=',\n', + mod=['!s', ':.4']) + return '{}(\n{}\n)'.format(self.__class__.__name__, indent(inner_str)) + class CupyTensorSpaceConstWeighting(ConstWeighting): From 8564f2fadc6496ae7b8a56ef62f3211f24d6b09f Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Fri, 24 Nov 2017 00:14:38 +0100 Subject: [PATCH 04/38] WIP: fix custom cupy kernels --- odl/space/cupy_tensors.py | 150 +++++++++++++++++++++++++------------- 1 file changed, 98 insertions(+), 52 deletions(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index 9577daecd50..988e7985354 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -22,7 +22,8 @@ Weighting, ArrayWeighting, ConstWeighting, CustomInner, CustomNorm, CustomDist) from odl.util import ( - array_str, dtype_str, is_floating_dtype, signature_string, indent) + array_str, dtype_str, is_floating_dtype, real_dtype, + signature_string, indent) try: import cupy @@ -1853,6 +1854,15 @@ def _weighting(weights, exponent): # Kernels for space functions +# +# T = generic type +# R = real floating point type, usually for norm or dist output +# W = real type for weights, can be floating point or integer +# +# Note: the kernels with an output type that does not also occur as an input +# type must be called with output argument since the output type cannot be +# inferred. The ouptut array must have shape `()` for full reduction. + if CUPY_AVAILABLE: dotw = cupy.ReductionKernel(in_params='T x, T y, W w', @@ -1864,7 +1874,7 @@ def _weighting(weights, exponent): name='dotw') nrm0 = cupy.ReductionKernel(in_params='T x', - out_params='int64 res', + out_params='R res', map_expr='x != 0', reduce_expr='a + b', post_map_expr='res = a', @@ -1872,7 +1882,7 @@ def _weighting(weights, exponent): name='nrm0') nrm1 = cupy.ReductionKernel(in_params='T x', - out_params='T res', + out_params='R res', map_expr='abs(x)', reduce_expr='a + b', post_map_expr='res = a', @@ -1880,7 +1890,7 @@ def _weighting(weights, exponent): name='nrm1w') nrm1w = cupy.ReductionKernel(in_params='T x, W w', - out_params='T res', + out_params='R res', map_expr='abs(x) * w', reduce_expr='a + b', post_map_expr='res = a', @@ -1888,7 +1898,7 @@ def _weighting(weights, exponent): name='nrm1w') nrm2 = cupy.ReductionKernel(in_params='T x', - out_params='T res', + out_params='R res', map_expr='x * x', reduce_expr='a + b', post_map_expr='res = sqrt(a)', @@ -1896,7 +1906,7 @@ def _weighting(weights, exponent): name='nrm2') nrm2w = cupy.ReductionKernel(in_params='T x, W w', - out_params='T res', + out_params='R res', map_expr='x * x * w', reduce_expr='a + b', post_map_expr='res = sqrt(a)', @@ -1904,7 +1914,7 @@ def _weighting(weights, exponent): name='nrm2w') nrminf = cupy.ReductionKernel(in_params='T x', - out_params='T res', + out_params='R res', map_expr='abs(x)', reduce_expr='a > b ? a : b', post_map_expr='res = a', @@ -1912,23 +1922,23 @@ def _weighting(weights, exponent): name='nrminf') nrmneginf = cupy.ReductionKernel(in_params='T x', - out_params='T res', + out_params='R res', map_expr='abs(x)', reduce_expr='a > b ? b : a', post_map_expr='res = a', identity='0', name='nrmneginf') - nrmp = cupy.ReductionKernel(in_params='T x, T p', - out_params='T res', + nrmp = cupy.ReductionKernel(in_params='T x, R p', + out_params='R res', map_expr='pow(abs(x), p)', reduce_expr='a + b', post_map_expr='res = pow(a, 1 / p)', identity='0', name='nrmp') - nrmpw = cupy.ReductionKernel(in_params='T x, T p, W w', - out_params='T res', + nrmpw = cupy.ReductionKernel(in_params='T x, R p, W w', + out_params='R res', map_expr='pow(abs(x), p) * w', reduce_expr='a + b', post_map_expr='res = pow(a, 1 / p)', @@ -1936,7 +1946,7 @@ def _weighting(weights, exponent): name='nrmpw') dist0 = cupy.ReductionKernel(in_params='T x, T y', - out_params='int64 res', + out_params='R res', map_expr='x != y', reduce_expr='a + b', post_map_expr='res = a', @@ -1944,7 +1954,7 @@ def _weighting(weights, exponent): name='dist0') dist1 = cupy.ReductionKernel(in_params='T x, T y', - out_params='T res', + out_params='R res', map_expr='abs(x - y)', reduce_expr='a + b', post_map_expr='res = a', @@ -1960,23 +1970,23 @@ def _weighting(weights, exponent): name='dist1w') dist2 = cupy.ReductionKernel(in_params='T x, T y', - out_params='T res', - map_expr='(x - y) * (x - y)', + out_params='R res', + map_expr='abs(x - y) * abs(x - y)', reduce_expr='a + b', post_map_expr='res = sqrt(a)', identity='0', name='dist2') dist2w = cupy.ReductionKernel(in_params='T x, T y, W w', - out_params='T res', - map_expr='(x - y) * (x - y) * w', + out_params='R res', + map_expr='abs(x - y) * abs(x - y) * w', reduce_expr='a + b', post_map_expr='res = sqrt(a)', identity='0', name='dist2w') distinf = cupy.ReductionKernel(in_params='T x, T y', - out_params='T res', + out_params='R res', map_expr='abs(x - y)', reduce_expr='a > b ? a : b', post_map_expr='res = a', @@ -1984,23 +1994,23 @@ def _weighting(weights, exponent): name='distinf') distneginf = cupy.ReductionKernel(in_params='T x, T y', - out_params='T res', + out_params='R res', map_expr='abs(x - y)', reduce_expr='a > b ? b : a', post_map_expr='res = a', identity='0', name='distneginf') - distp = cupy.ReductionKernel(in_params='T x, T y, T p', - out_params='T res', + distp = cupy.ReductionKernel(in_params='T x, T y, R p', + out_params='R res', map_expr='pow(abs(x - y), p)', reduce_expr='a + b', post_map_expr='res = pow(a, 1 / p)', identity='0', name='distp') - distpw = cupy.ReductionKernel(in_params='T x, T y, T p, W w', - out_params='T res', + distpw = cupy.ReductionKernel(in_params='T x, T y, R p, W w', + out_params='R res', map_expr='pow(abs(x - y), p) * w', reduce_expr='a + b', post_map_expr='res = pow(a, 1 / p)', @@ -2066,18 +2076,26 @@ def norm(self, x): norm : float The norm of the provided tensor. """ + # Define scalar output array + if is_floating_dtype(x.dtype): + out_dtype = real_dtype(x.dtype) + else: + out_dtype = float + out = cupy.empty((), dtype=out_dtype) + + # Run reduction kernel (returns the output) if self.exponent == 0: - return float(nrm0(x.data)) + return float(nrm0(x.data, out)) elif self.exponent == 1: - return float(nrm1w(x.data, self.array)) + return float(nrm1w(x.data, self.array, out)) elif self.exponent == 2: - return float(nrm2w(x.data, self.array)) + return float(nrm2w(x.data, self.array, out)) elif self.exponent == float('inf'): - return float(nrminf(x.data)) + return float(nrminf(x.data, out)) elif self.exponent == -float('inf'): - return float(nrmneginf(x.data)) + return float(nrmneginf(x.data, out)) else: - return float(nrmpw(x.data, self.exponent, self.array)) + return float(nrmpw(x.data, self.exponent, self.array, out)) def dist(self, x1, x2): """Return the weighted distance of two tensors. @@ -2092,18 +2110,27 @@ def dist(self, x1, x2): dist : float The distance between the provided tensors. """ + # Define scalar output array + if is_floating_dtype(x1.dtype): + out_dtype = real_dtype(x1.dtype) + else: + out_dtype = float + out = cupy.empty((), dtype=out_dtype) + + # Run reduction kernel (returns the output) if self.exponent == 0: - return float(dist0(x1.data, x2.data)) + return float(dist0(x1.data, x2.data, out)) elif self.exponent == 1: - return float(dist1w(x1.data, x2.data, self.array)) + return float(dist1w(x1.data, x2.data, self.array, out)) elif self.exponent == 2: - return float(dist2w(x1.data, x2.data, self.array)) + return float(dist2w(x1.data, x2.data, self.array, out)) elif self.exponent == float('inf'): - return float(distinf(x1.data, x2.data, self.array)) + return float(distinf(x1.data, x2.data, out)) elif self.exponent == -float('inf'): - return float(distneginf(x1.data, x2.data, self.array)) + return float(distneginf(x1.data, x2.data, out)) else: - return float(distpw(x1.data, x2.data, self.exponent, self.array)) + return float(distpw(x1.data, x2.data, self.exponent, self.array, + out)) # TODO: remove repr_part and __repr__ when cupy.ndarray.__array__ # is implemented. See @@ -2111,8 +2138,9 @@ def dist(self, x1, x2): @property def repr_part(self): """String usable in a space's ``__repr__`` method.""" - # TODO: use edgeitems - arr_str = array_str(cupy.asnumpy(self.array), nprint=10) + maxsize_full_print = 2 * np.get_printoptions()['edgeitems'] + arr_str = array_str(cupy.asnumpy(self.array), + nprint=maxsize_full_print) optargs = [('weighting', arr_str, ''), ('exponent', self.exponent, 2.0)] return signature_string([], optargs, sep=',\n', @@ -2120,8 +2148,10 @@ def repr_part(self): def __repr__(self): """Return ``repr(self)``.""" - # TODO: use edgeitems - posargs = [array_str(cupy.asnumpy(self.array), nprint=10)] + maxsize_full_print = 2 * np.get_printoptions()['edgeitems'] + arr_str = array_str(cupy.asnumpy(self.array), + nprint=maxsize_full_print) + posargs = [arr_str] optargs = [('exponent', self.exponent, 2.0)] inner_str = signature_string(posargs, optargs, sep=',\n', mod=['!s', ':.4']) @@ -2183,10 +2213,18 @@ def norm(self, x): norm : float The norm of the tensor. """ + # Define scalar output array + if is_floating_dtype(x.dtype): + out_dtype = real_dtype(x.dtype) + else: + out_dtype = float + out = cupy.empty((), dtype=out_dtype) + + # Run reduction kernel (returns the output) if self.exponent == 0: - return float(nrm0(x.data)) + return float(nrm0(x.data, out)) elif self.exponent == 1: - return float(self.const * nrm1(x.data)) + return float(self.const * nrm1(x.data, out)) elif self.exponent == 2: # We try to use cuBLAS nrm2 try: @@ -2208,14 +2246,14 @@ def norm(self, x): return float(np.sqrt(self.const) * norm) # Cannot use cuBLAS, fall back to custom kernel - return float(np.sqrt(self.const) * nrm2(x.data)) + return float(np.sqrt(self.const) * nrm2(x.data, out)) elif self.exponent == float('inf'): - return float(nrminf(x.data)) + return float(nrminf(x.data, out)) elif self.exponent == -float('inf'): - return float(nrmneginf(x.data)) + return float(nrmneginf(x.data, out)) else: return float(self.const ** (1 / self.exponent) * - nrmp(x, self.exponent)) + nrmp(x, self.exponent, out)) def dist(self, x1, x2): """Return the weighted distance between two tensors. @@ -2230,21 +2268,29 @@ def dist(self, x1, x2): dist : float The distance between the tensors. """ + # Define scalar output array + if is_floating_dtype(x1.dtype): + out_dtype = real_dtype(x1.dtype) + else: + out_dtype = float + out = cupy.empty((), dtype=out_dtype) + + # Run reduction kernel (returns the output) if self.exponent == 0: - return float(dist0(x1.data, x2.data)) + return float(dist0(x1.data, x2.data, out)) elif self.exponent == 1: - return float(self.const * dist1(x1.data, x2.data)) + return float(self.const * dist1(x1.data, x2.data, out)) elif self.exponent == 2: # cuBLAS nrm2(x1 - x2) would probably be faster, but would # require a copy, so we don't do it - return float(np.sqrt(self.const) * dist2(x1.data, x2.data)) + return float(np.sqrt(self.const) * dist2(x1.data, x2.data, out)) elif self.exponent == float('inf'): - return float(distinf(x1.data, x2.data)) + return float(distinf(x1.data, x2.data, out)) elif self.exponent == -float('inf'): - return float(distneginf(x1.data, x2.data)) + return float(distneginf(x1.data, x2.data, out)) else: return float(self.const ** (1 / self.exponent) * - distp(x1.data, x2.data, self.exponent)) + distp(x1.data, x2.data, self.exponent, out)) class CupyTensorSpaceCustomInner(CustomInner): From f2085f8156da73cbe4062add4550aa8d8b6fc3aa Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Fri, 24 Nov 2017 00:15:02 +0100 Subject: [PATCH 05/38] MAINT: minor fixes --- odl/space/npy_tensors.py | 3 +-- odl/util/testutils.py | 2 +- odl/util/utility.py | 2 +- 3 files changed, 3 insertions(+), 4 deletions(-) diff --git a/odl/space/npy_tensors.py b/odl/space/npy_tensors.py index af0c91c156a..757351d465f 100644 --- a/odl/space/npy_tensors.py +++ b/odl/space/npy_tensors.py @@ -220,8 +220,7 @@ def __init__(self, shape, dtype=None, **kwargs): """ super(NumpyTensorSpace, self).__init__(shape, dtype) if self.dtype.char not in self.available_dtypes(): - raise ValueError('`dtype` {!r} not supported' - ''.format(dtype_str(dtype))) + raise ValueError('`dtype` {!r} not supported'.format(dtype)) dist = kwargs.pop('dist', None) norm = kwargs.pop('norm', None) diff --git a/odl/util/testutils.py b/odl/util/testutils.py index a0f54f76452..16ece42b31d 100644 --- a/odl/util/testutils.py +++ b/odl/util/testutils.py @@ -305,7 +305,7 @@ def noise_array(space): Returns ------- - noise_array : `numpy.ndarray` element + noise_array : `numpy.ndarray` Array with white noise such that ``space.element``'s can be created from it. diff --git a/odl/util/utility.py b/odl/util/utility.py index 828ee0ed5d1..ce23fe14dc8 100644 --- a/odl/util/utility.py +++ b/odl/util/utility.py @@ -438,7 +438,7 @@ def real_dtype(dtype, default=None): """ dtype, dtype_in = np.dtype(dtype), dtype - if is_real_floating_dtype(dtype): + if is_real_dtype(dtype): return dtype try: From 0654031e4b3dd250f36ab9bd8e9ce003090e1083 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Fri, 24 Nov 2017 00:15:33 +0100 Subject: [PATCH 06/38] WIP: implement cupy space unit tests --- odl/test/space/tensors_test.py | 121 +++++++++++++++++++++++++++++---- 1 file changed, 106 insertions(+), 15 deletions(-) diff --git a/odl/test/space/tensors_test.py b/odl/test/space/tensors_test.py index 1316f850c6f..24bc061ba20 100644 --- a/odl/test/space/tensors_test.py +++ b/odl/test/space/tensors_test.py @@ -22,11 +22,20 @@ NumpyTensorSpaceConstWeighting, NumpyTensorSpaceArrayWeighting, NumpyTensorSpaceCustomInner, NumpyTensorSpaceCustomNorm, NumpyTensorSpaceCustomDist) +from odl.space.cupy_tensors import ( + CupyTensor, CupyTensorSpace, + CupyTensorSpaceConstWeighting, CupyTensorSpaceArrayWeighting, + CupyTensorSpaceCustomInner, CupyTensorSpaceCustomNorm, + CupyTensorSpaceCustomDist, + CUPY_AVAILABLE) from odl.util.testutils import ( all_almost_equal, all_equal, simple_fixture, noise_array, noise_element, noise_elements) from odl.util.ufuncs import UFUNCS +if CUPY_AVAILABLE: + import cupy + # --- Test helpers --- # @@ -39,14 +48,22 @@ # when a new impl is available. def _pos_array(space): - """Create an array with positive real entries in ``space``.""" - return np.abs(noise_array(space)) + 0.1 + """Create a Numpy array with positive real entries for ``space``.""" + arr = np.abs(noise_array(space)) + 0.1 + if space.impl == 'numpy': + return arr + elif space.impl == 'cupy': + return cupy.asarray(arr) + else: + assert False def _array_cls(impl): """Return the array class for given impl.""" if impl == 'numpy': return np.ndarray + elif impl == 'cupy': + return cupy.ndarray else: assert False @@ -55,6 +72,8 @@ def _odl_tensor_cls(impl): """Return the ODL tensor class for given impl.""" if impl == 'numpy': return NumpyTensor + elif impl == 'cupy': + return CupyTensor else: assert False @@ -74,6 +93,21 @@ def _weighting_cls(impl, kind): return NumpyTensorSpaceCustomDist else: assert False + + elif impl == 'cupy': + if kind == 'array': + return CupyTensorSpaceArrayWeighting + elif kind == 'const': + return CupyTensorSpaceConstWeighting + elif kind == 'inner': + return CupyTensorSpaceCustomInner + elif kind == 'norm': + return CupyTensorSpaceCustomNorm + elif kind == 'dist': + return CupyTensorSpaceCustomDist + else: + assert False + else: assert False @@ -121,6 +155,9 @@ def test_init_npy_tspace(): NumpyTensorSpace((3, 4), dtype=complex, exponent=float('inf')) NumpyTensorSpace((3, 4), dtype='S1') + with pytest.raises(ValueError): + NumpyTensorSpace((3, 4), dtype=object) + # Alternative constructor odl.tensor_space((3, 4)) odl.tensor_space((3, 4), dtype=int) @@ -160,19 +197,73 @@ def test_init_npy_tspace(): odl.rn((3, 4), weighting=weight_arr) -def test_init_tspace_weighting(weight, exponent, odl_tspace_impl): - """Test if weightings during init give the correct weighting classes.""" - impl = odl_tspace_impl - space = odl.tensor_space((3, 4), weighting=weight, exponent=exponent, - impl=impl) +def test_init_cupy_tspace(): + """Test initialization patterns and options for ``CupyTensorSpace``.""" + if not CUPY_AVAILABLE: + pytest.skip('cupy backend not available') - if impl == 'numpy': - if isinstance(weight, np.ndarray): - weighting_cls = _weighting_cls(impl, 'array') - else: - weighting_cls = _weighting_cls(impl, 'const') + # Basic class constructor + CupyTensorSpace((3, 4)) + CupyTensorSpace((3, 4), dtype=int) + CupyTensorSpace((3, 4), dtype=float) + CupyTensorSpace((3, 4), dtype=complex) + CupyTensorSpace((3, 4), dtype=complex, exponent=1.0) + CupyTensorSpace((3, 4), dtype=complex, exponent=float('inf')) + + with pytest.raises(ValueError): + CupyTensorSpace((3, 4), dtype='S1') + with pytest.raises(ValueError): + CupyTensorSpace((3, 4), dtype=object) + + # Alternative constructor + odl.tensor_space((3, 4), impl='cupy') + odl.tensor_space((3, 4), dtype=int, impl='cupy') + odl.tensor_space((3, 4), exponent=1.0, impl='cupy') + + # Constructors for real spaces + odl.rn((3, 4), impl='cupy') + odl.rn((3, 4), dtype='float32', impl='cupy') + odl.rn(3, impl='cupy') + odl.rn(3, dtype='float32', impl='cupy') + + # Works only for real data types + with pytest.raises(ValueError): + odl.rn((3, 4), complex, impl='cupy') + with pytest.raises(ValueError): + odl.rn(3, int, impl='cupy') + with pytest.raises(ValueError): + odl.rn(3, 'S1', impl='cupy') + + # Constructors for complex spaces + odl.cn((3, 4), impl='cupy') + odl.cn((3, 4), dtype='complex64', impl='cupy') + odl.cn(3, impl='cupy') + odl.cn(3, dtype='complex64', impl='cupy') + + # Works only for complex data types + with pytest.raises(ValueError): + odl.cn((3, 4), float, impl='cupy') + with pytest.raises(ValueError): + odl.cn(3, 'S1', impl='cupy') + + # Init with weights or custom space functions + weight_const = 1.5 + weight_arr = _pos_array(odl.rn((3, 4), float)) + + odl.rn((3, 4), weighting=weight_const, impl='cupy') + odl.rn((3, 4), weighting=weight_arr, impl='cupy') + + +def test_init_tspace_weighting(weight, exponent, tspace_impl): + """Test if weightings during init give the correct weighting classes.""" + if tspace_impl == 'cupy' and isinstance(weight, np.ndarray): + # Need cast before using in space creation since + # ArrayWeighting.__eq__ uses `arr1 is arr2` to check arrays + weight = cupy.asarray(weight) + elif isinstance(weight, _array_cls(tspace_impl)): + weighting_cls = _weighting_cls(tspace_impl, 'array') else: - assert False + weighting_cls = _weighting_cls(tspace_impl, 'const') weighting = weighting_cls(weight, exponent) @@ -193,8 +284,8 @@ def test_init_tspace_weighting(weight, exponent, odl_tspace_impl): bad_dtype = np.ones((3, 4), dtype=complex) odl.tensor_space((3, 4), weighting=bad_dtype) - with pytest.raises(TypeError): - odl.tensor_space((3, 4), weighting=1j) # float() conversion + with pytest.raises(TypeError): + odl.tensor_space((3, 4), weighting=1j) # float() conversion def test_properties(odl_tspace_impl): From dab3926fd1b451407d117f2ad0f5e80aa4fb1704 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Fri, 24 Nov 2017 15:36:29 +0100 Subject: [PATCH 07/38] WIP: fix cublas stuff --- odl/space/cupy_tensors.py | 187 +++++++++++++++++++++++++++++++++++--- 1 file changed, 174 insertions(+), 13 deletions(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index 988e7985354..c5aee3b3cd7 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -51,24 +51,176 @@ name='lico') -def fallback_scal(a, x): +def _fallback_scal(a, x): + """Fallback implementation of ``scal`` when cuBLAS is not applicable.""" x *= a return x -def fallback_axpy(a, x, y): +def _fallback_axpy(a, x, y): + """Fallback implementation of ``axpy`` when cuBLAS is not applicable.""" return lico(a, x, 1, y, y) -def _flat_inc(arr): - """Compute the flat element stride for cuBLAS if possible, else raise.""" - flat_inc = min(arr.strides) // arr.itemsize - stride = min(arr.strides) - for n, s in zip(sorted(arr.shape)[:-1], sorted(arr.strides)[1:]): - next_stride = stride * n - if s != next_stride: +def _get_flat_inc(arr1, arr2=None): + """Return flat index increment(s) for cuBLAS, raise if not applicable. + + This function checks if the array stride(s) allow usage of cuBLAS + and returns the ``incx`` (and ``incy``) parameters needed for the + cuBLAS functions. If the array strides ar such that cuBLAS cannot + be applied, a ``ValueError`` is raised, triggering a fallback + implementation. + + For **1 array**, the conditions to be fulfilled are + + - the strides do not contain 0 and + - the memory of the array has constant stride. + + The second point applies to + + - contiguous arrays, + - chunks of arrays along the **slowest-varying axis** (``arr[1:5, ...]`` + for C-contiguous ``arr``), and + - strided slices along the **fastest-varying axis** (``arr[..., ::2]`` + for C-contiguous ``arr``). + + For **2 arrays**, both arrays must + + - fulfill the "1 array" conditions individually, + - have the same total size, and + - the axis order of both must be the same in the sense that the same + index array sorts the strides of both arrays in ascending order. + + Parameters + ---------- + arr1 : cupy.core.core.ndarray + Array to check for compatibility. + arr2 : cupy.core.core.ndarray, optional + Second array to check for compatibility, by itself and with ``arr1``. + + Returns + ------- + flat_inc1 : int + Memory stride (in terms of elements, not bytes) of ``arr1``. + flat_inc2 : int or None + Memory stride of ``arr2`` if provided, otherwise ``None``. + + Raises + ------ + ValueError + If the conditions for cuBLAS compatibility are not met. + + Examples + -------- + >>> arr_c = cupy.zeros((4, 4, 4)) + >>> arr_c.strides + (128, 32, 8) + >>> arr_f = cupy.asfortranarray(arr_c) + >>> arr_f.strides + (8, 32, 128) + + Contiguous arrays are compatible by and with themselves, but not with + arrays that are contiguous in a different ordering: + + >>> _get_flat_inc(arr_c) + 1 + >>> _get_flat_inc(arr_c, arr_c) + (1, 1) + True + >>> _get_flat_inc(arr_f) + 1 + >>> _get_flat_inc(arr_f, arr_f) + (1, 1) + >>> _get_flat_inc(arr_c, arr_f) + Traceback (most recent call last): + ... + ValueError + + Slicing in the **fastest** axis is allowed as it results in a constant + stride in the flat memory. Slicing with stride in any other axis is + results in incompatibility. + + C ordering (last axis fastest): + + >>> half_arr_0_c = cupy.zeros((2, 4, 4)) + >>> half_arr_1_c = cupy.zeros((4, 2, 4)) + >>> half_arr_2_c = cupy.zeros((4, 4, 2)) + >>> _get_flat_inc(arr_c[::2, :, :], half_arr_0_c) + Traceback (most recent call last): + ... + ValueError + >>> _get_flat_inc(arr_c[:, ::2, :], half_arr_1_c) + Traceback (most recent call last): + ... + ValueError + >>> _get_flat_inc(arr_c[:, :, ::2], half_arr_2_c) + (2, 1) + + Fortran ordering (first axis fastest): + + >>> half_arr_0_f = cupy.asfortranarray(half_arr_0_c) + >>> half_arr_1_f = cupy.asfortranarray(half_arr_1_c) + >>> half_arr_2_f = cupy.asfortranarray(half_arr_2_c) + >>> _get_flat_inc(arr_f[::2, :, :], half_arr_0_f) + (2, 1) + >>> _get_flat_inc(arr_f[:, ::2, :], half_arr_1_f) + Traceback (most recent call last): + ... + ValueError + >>> _get_flat_inc(arr_f[:, :, ::2], half_arr_2_f) + Traceback (most recent call last): + ... + ValueError + + Axes swapped (middle axis fastest): + + >>> arr_s = cupy.swapaxes(arr_c, 1, 2) + >>> arr_s.strides + (128, 8, 32) + >>> half_arr_0_s = cupy.swapaxes(half_arr_0_c, 1, 2) + >>> half_arr_1_s = cupy.swapaxes(half_arr_1_c, 1, 2) + >>> half_arr_2_s = cupy.swapaxes(half_arr_2_c, 1, 2) + >>> _get_flat_inc(arr_s[:, :, ::2], half_arr_0_s) + Traceback (most recent call last): + ... + ValueError + >>> _get_flat_inc(arr_s[:, ::2, :], half_arr_1_s) + (2, 1) + >>> _get_flat_inc(arr_s[:, :, ::2], half_arr_2_s) + Traceback (most recent call last): + ... + ValueError + """ + # Zero strides not allowed + if 0 in arr1.strides: + raise ValueError + if arr2 is not None and 0 in arr2.strides: + raise ValueError + + # Candidate for flat_inc of array 1 + arr1_flat_inc = min(arr1.strides) // arr1.itemsize + + # Check if the strides are as in a contiguous array (after reordering + # the axes), except for the fastest axis. We allow arbitrary axis order + # for operations on the whole array at once, as long as it can be + # indexed with a single flat index and stride. + arr1_ax_order = np.argsort(arr1.strides) # ascending + arr1_sorted_shape = np.take(arr1.shape, arr1_ax_order) + arr1_sorted_shape[0] *= arr1_flat_inc + arr1_elem_strides = np.take(arr1.strides, arr1_ax_order) // arr1.itemsize + if np.any(np.cumprod(arr1_sorted_shape[:-1]) != arr1_elem_strides[1:]): + raise ValueError + + if arr2 is None: + return arr1_flat_inc + else: + arr2_flat_inc = _get_flat_inc(arr2) + if arr1.size != arr2.size: raise ValueError - return flat_inc + if np.any(np.diff(np.take(arr2.strides, arr1_ax_order)) < 0): + # Strides of arr2 are not sorted by axis order of arr1 + raise ValueError + return arr1_flat_inc, arr2_flat_inc def _cublas_func(name, dtype): @@ -80,11 +232,11 @@ def _cublas_func(name, dtype): Raw function name without prefix, e.g., ``'axpy'``. dtype : Numpy dtype specifier for which the cuBLAS function should be - used. Must be either single or double precision float. + used. Must be single or double precision float or complex. Raises ------ - ValueError : + ValueError If the data type is not supported by cuBLAS. """ dtype, dtype_in = np.dtype(dtype), dtype @@ -103,7 +255,16 @@ def _cublas_func(name, dtype): def _get_scal_axpy(x1, x2): - """Return implementations of scal and axpy suitable for the inputs.""" + """Return implementations of scal and axpy suitable for the inputs. + + If the inputs are suitable, a cuBLAS implementation is returned, otherwise + a fallback implementation. To be suitable for cuBLAS, both arrays + must + + - have single or double precision float or complex data type and + - have a single integer stride when flattened, i.e., be contiguous + or (this allows using) + """ try: incx1 = _flat_inc(x1.data) incx2 = _flat_inc(x2.data) From ffde709cb901dfd03ed327db50212bd5d0ce9e58 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Sat, 25 Nov 2017 12:33:15 +0100 Subject: [PATCH 08/38] WIP: fix cublas, weight propagation, implement (add,multiply).(reduce,accumulate) --- odl/space/cupy_tensors.py | 241 ++++++++++++++++++++++++++++---------- odl/space/npy_tensors.py | 13 +- odl/util/testutils.py | 6 +- 3 files changed, 192 insertions(+), 68 deletions(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index c5aee3b3cd7..69d5f5c4810 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -22,7 +22,7 @@ Weighting, ArrayWeighting, ConstWeighting, CustomInner, CustomNorm, CustomDist) from odl.util import ( - array_str, dtype_str, is_floating_dtype, real_dtype, + array_str, dtype_str, is_floating_dtype, is_numeric_dtype, real_dtype, signature_string, indent) try: @@ -45,8 +45,8 @@ if CUPY_AVAILABLE: - lico = cupy.ElementwiseKernel(in_params='T a, T x, T b, T y', - out_params='T z', + lico = cupy.ElementwiseKernel(in_params='X a, X x, Y b, Y y', + out_params='Z z', operation='z = a * x + b * y;', name='lico') @@ -102,8 +102,8 @@ def _get_flat_inc(arr1, arr2=None): ------- flat_inc1 : int Memory stride (in terms of elements, not bytes) of ``arr1``. - flat_inc2 : int or None - Memory stride of ``arr2`` if provided, otherwise ``None``. + flat_inc2 : int or None, optional + Memory stride of ``arr2``. Returned only if ``arr2`` was provided. Raises ------ @@ -126,7 +126,6 @@ def _get_flat_inc(arr1, arr2=None): 1 >>> _get_flat_inc(arr_c, arr_c) (1, 1) - True >>> _get_flat_inc(arr_f) 1 >>> _get_flat_inc(arr_f, arr_f) @@ -254,51 +253,113 @@ def _cublas_func(name, dtype): return getattr(cupy.cuda.cublas, prefix + name) -def _get_scal_axpy(x1, x2): - """Return implementations of scal and axpy suitable for the inputs. +def _get_scal(arr): + """Return a ``scal`` implementation suitable for the input. - If the inputs are suitable, a cuBLAS implementation is returned, otherwise - a fallback implementation. To be suitable for cuBLAS, both arrays - must + If possible, a cuBLAS implementation is returned, otherwise a fallback. - - have single or double precision float or complex data type and - - have a single integer stride when flattened, i.e., be contiguous - or (this allows using) + In general, cuBLAS requires single or double precision float or complex + data type, and operates on linear memory with constant stride. + The latter is the case for + + - contiguous arrays, + - chunks of arrays along the **slowest-varying axis** (``arr[1:5, ...]`` + for C-contiguous ``arr``), and + - strided slices along the **fastest-varying axis** (``arr[..., ::2]`` + for C-contiguous ``arr``). + + Note that CuPy's cuBLAS wrapper does not necessarily implement all + functions for all four data types. See + https://github.com/cupy/cupy/blob/master/cupy/cuda/cublas.pyx + for details. + + Parameters + ---------- + arr : cupy.core.core.ndarray + Array to which ``scal`` should be applied. + + Returns + ------- + scal : callable + Implementation of ``x <- a * x``. """ try: - incx1 = _flat_inc(x1.data) - incx2 = _flat_inc(x2.data) + incx = _get_flat_inc(arr) except ValueError: - use_cublas = False - else: - use_cublas = True + return _fallback_scal - if use_cublas: - try: - scal_cublas = _cublas_func('scal', x1.dtype) - except (ValueError, AttributeError): - scal = fallback_scal - else: - def scal(a, x): - with cupy.cuda.Device(x1.device) as dev: - return scal_cublas( - dev.cublas_handle, x.data.size, a, x.data.ptr, incx1) + try: + _cublas_scal = _cublas_func('scal', arr.dtype) + except (ValueError, AttributeError): + return _fallback_scal - try: - axpy_cublas = _cublas_func('axpy', x1.dtype) - except (ValueError, AttributeError): - axpy = fallback_axpy - else: - def axpy(a, x, y): - with cupy.cuda.Device(x1.device) as dev: - return axpy_cublas( - dev.cublas_handle, x.data.size, a, - x.data.ptr, incx1, y.data.ptr, incx2) - else: - scal = fallback_scal - axpy = fallback_axpy + def scal(a, x): + """Implement ``x <- a * x`` with constant ``a``.""" + return _cublas_scal( + x.data.device.cublas_handle, x.size, a, x.data.ptr, incx) - return scal, axpy + scal.__name__ = scal.__qualname__ = '_cublas_scal' + return scal + + +def _get_axpy(arr1, arr2): + """Return an ``axpy`` implementation suitable for the inputs. + + If possible, a cuBLAS implementation is returned, otherwise a fallback. + + In general, cuBLAS requires single or double precision float or complex + data type, and operates on linear memory with constant stride. + The latter is the case for + + - contiguous arrays, + - chunks of arrays along the **slowest-varying axis** (``arr[1:5, ...]`` + for C-contiguous ``arr``), and + - strided slices along the **fastest-varying axis** (``arr[..., ::2]`` + for C-contiguous ``arr``). + + Furthermore, the two arrays must + + - be on the same device, + - have the same total size, and + - the axis order of both must be the same in the sense that the same + index array sorts the strides of both arrays in ascending order. + + Note that CuPy's cuBLAS wrapper does not necessarily implement all + functions for all four data types. See + https://github.com/cupy/cupy/blob/master/cupy/cuda/cublas.pyx + for details. + + Parameters + ---------- + arr1, arr2 : cupy.core.core.ndarray + Arrays to which ``axpy`` should be applied. + + Returns + ------- + axpy : callable + Implementation of ``y <- a * x + y``. + """ + if arr1.dtype != arr2.dtype or arr1.device != arr2.device: + return _fallback_axpy + + try: + incx1, incx2 = _get_flat_inc(arr1, arr2) + except ValueError: + return _fallback_axpy + + try: + _cublas_axpy = _cublas_func('axpy', arr1.dtype) + except (ValueError, AttributeError): + return _fallback_axpy + + def axpy(a, x, y): + """Implement ``y <- a * x + y`` with constant ``a``.""" + return _cublas_axpy( + x.data.device.cublas_handle, x.size, a, x.data.ptr, incx1, + y.data.ptr, incx2) + + axpy.__name__ = axpy.__qualname__ = '_cublas_axpy' + return axpy def _lincomb_impl(a, x1, b, x2, out): @@ -307,7 +368,9 @@ def _lincomb_impl(a, x1, b, x2, out): This implementation is a highly optimized, considering all special cases of array alignment and special scalar values 0 and 1 separately. """ - scal, axpy = _get_scal_axpy(x1, x2) + scal1 = _get_scal(x1.data) + scal2 = _get_scal(x2.data) + axpy = _get_axpy(x1.data, x2.data) if a == 0 and b == 0: # out <- 0 @@ -320,7 +383,7 @@ def _lincomb_impl(a, x1, b, x2, out): if b == 1: pass else: - scal(b, out.data) + scal2(b, out.data) else: # out <- b * x2 if b == 1: @@ -335,7 +398,7 @@ def _lincomb_impl(a, x1, b, x2, out): if a == 1: pass else: - scal(a, out.data) + scal1(a, out.data) else: # out <- a * x1 if a == 1: @@ -354,7 +417,7 @@ def _lincomb_impl(a, x1, b, x2, out): elif a + b == 1: pass else: - scal(a + b, out.data) + scal1(a + b, out.data) elif out is x1 and a == 1: # out <-- out + b * x2 axpy(b, x2.data, out.data) @@ -1284,8 +1347,20 @@ def __getitem__(self, indices): raise RuntimeError("no conversion for dtype {}" "".format(arr.dtype)) else: + if is_numeric_dtype(self.dtype): + weighting = self.space.weighting + else: + weighting = None + + if isinstance(weighting, CupyTensorSpaceArrayWeighting): + weighting = weighting.array[indices] + elif isinstance(weighting, CupyTensorSpaceConstWeighting): + # Axes were removed, cannot infer new constant + if arr.ndim != self.ndim: + weighting = None + space = type(self.space)(arr.shape, dtype=self.dtype, - device=self.device) + device=self.device, weighting=weighting) return space.element(arr) def __setitem__(self, indices, values): @@ -1657,8 +1732,13 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): use_native = False else: native_ufunc = getattr(cupy, ufunc.__name__, None) - use_native = (native_ufunc is not None and - hasattr(native_ufunc, method)) + # Manual assignment for sum, cumsum, prod and cumprod + if (ufunc in (np.add, np.multiply) and + method in ('reduce', 'accumulate')): + use_native = native_ufunc is not None + else: + use_native = (native_ufunc is not None and + hasattr(native_ufunc, method)) # Assign to `out` or `out1` and `out2`, respectively, unwrapping the # data container @@ -1687,8 +1767,10 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): if use_native: # TODO: remove when upstream issue is fixed - # For native ufuncs, we turn non-scalar inputs into cupy arrays, - # as a workaround for https://github.com/cupy/cupy/issues/594 + # For native ufuncs, we turn non-scalar inputs into cupy arrays + # since cupy ufuncs do not accept array-like input. See + # https://github.com/cupy/cupy/issues/594 and + # https://github.com/odlgroup/odl/issues/1248 inputs, orig_inputs = [], inputs for inp in orig_inputs: if (isinstance(inp, cupy.ndarray) or @@ -1697,12 +1779,13 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): inputs.append(inp) else: inputs.append(cupy.array(inp)) - elif method != 'at': + elif not use_native and method != 'at': # TODO: remove when upstream issue is fixed # For non-native ufuncs (except `at`), we need ot cast our tensors # and Cupy arrays to Numpy arrays explicitly, since `__array__` # and friends are not implemented. See - # https://github.com/cupy/cupy/issues/589 + # https://github.com/cupy/cupy/issues/589 and + # https://github.com/odlgroup/odl/issues/1248 inputs, orig_inputs = [], inputs for inp in orig_inputs: if isinstance(inp, cupy.ndarray): @@ -1791,8 +1874,10 @@ def eval_at_via_npy(*inputs, **kwargs): new_inputs = (npy_arr,) + inputs[1:] super(CupyTensor, self).__array_ufunc__( ufunc, method, *new_inputs, **kwargs) - # Workaround for https://github.com/cupy/cupy/issues/593 - # TODO: use cupy_arr[:] = npy_arr when available + # Workaround for assignment cupy_arr[:] = npy_arr not + # working. See + # https://github.com/cupy/cupy/issues/593 and + # https://github.com/odlgroup/odl/issues/1248 cupy_arr.data.copy_from_host( npy_arr.ctypes.data_as(ctypes.c_void_p), npy_arr.nbytes) @@ -1810,6 +1895,41 @@ def eval_at_via_npy(*inputs, **kwargs): else: eval_at_via_npy(*inputs, **kwargs) + elif (ufunc in (np.add, np.multiply) and + method in ('reduce', 'accumulate')): + # These cases are implemented but not available as methods + # of cupy.ufunc. We do the manual assignment + if ufunc == np.add: + function = cupy.sum if method == 'reduce' else cupy.cumsum + else: + function = cupy.prod if method == 'reduce' else cupy.cumprod + + res = function(*inputs, **kwargs) + + # Shortcut for scalar or no return value + if np.isscalar(res): + # Happens for `reduce` with all axes + return res + + # Wrap result if necessary (lazily) + if out is None: + if is_floating_dtype(res.dtype): + if res.shape != self.shape: + # Don't propagate weighting if shape changes + weighting = CupyTensorSpaceConstWeighting( + 1.0, exponent) + spc_kwargs = {'weighting': weighting} + else: + spc_kwargs = {} + + res_space = type(self.space)( + res.shape, res.dtype, self.device, **spc_kwargs) + result = res_space.element(res) + else: + result = out_tuple[0] + + return result + else: # method != '__call__' kwargs['out'] = (out,) native_method = getattr(native_ufunc, method, None) @@ -2389,7 +2509,7 @@ def norm(self, x): elif self.exponent == 2: # We try to use cuBLAS nrm2 try: - incx = _flat_inc(x.data) + incx = _get_flat_inc(x.data) except ValueError: use_cublas = False else: @@ -2397,13 +2517,12 @@ def norm(self, x): if use_cublas: try: - nrm2_cublas = _cublas_func('nrm2', x.dtype) + _cublas_nrm2 = _cublas_func('nrm2', x.dtype) except (ValueError, AttributeError): pass else: - with cupy.cuda.Device(x.device) as dev: - norm = nrm2_cublas( - dev.cublas_handle, x.size, x.data_ptr, incx) + norm = _cublas_nrm2( + x.data.device.cublas_handle, x.size, x.data_ptr, incx) return float(np.sqrt(self.const) * norm) # Cannot use cuBLAS, fall back to custom kernel diff --git a/odl/space/npy_tensors.py b/odl/space/npy_tensors.py index 757351d465f..f0e04bc7795 100644 --- a/odl/space/npy_tensors.py +++ b/odl/space/npy_tensors.py @@ -1108,6 +1108,14 @@ def __getitem__(self, indices): weighting = self.space.weighting else: weighting = None + + if isinstance(weighting, NumpyTensorSpaceArrayWeighting): + weighting = weighting.array[indices] + elif isinstance(weighting, NumpyTensorSpaceConstWeighting): + # Axes were removed, cannot infer new constant + if arr.ndim != self.ndim: + weighting = None + space = type(self.space)( arr.shape, dtype=self.dtype, exponent=self.space.exponent, weighting=weighting) @@ -1124,11 +1132,6 @@ def __setitem__(self, indices, values): values : scalar, array-like or `NumpyTensor` The value(s) that are to be assigned. - If ``index`` is an integer, ``value`` must be a scalar. - - If ``index`` is a slice or a sequence of slices, ``value`` - must be broadcastable to the shape of the slice. - Examples -------- For 1d spaces, entries can be set with scalars or sequences of diff --git a/odl/util/testutils.py b/odl/util/testutils.py index 16ece42b31d..9339b65e081 100644 --- a/odl/util/testutils.py +++ b/odl/util/testutils.py @@ -687,12 +687,14 @@ def run_doctests(skip_if=False, **kwargs): Extra keyword arguments passed on to the ``doctest.testmod`` function. """ - from doctest import testmod, NORMALIZE_WHITESPACE, SKIP + from doctest import ( + testmod, NORMALIZE_WHITESPACE, SKIP, IGNORE_EXCEPTION_DETAIL) from pkg_resources import parse_version import odl import numpy as np - optionflags = kwargs.pop('optionflags', NORMALIZE_WHITESPACE) + optionflags = kwargs.pop('optionflags', + NORMALIZE_WHITESPACE | IGNORE_EXCEPTION_DETAIL) if skip_if: optionflags |= SKIP From 3298d74c8f9cea85aac1b46d125dadd30ee4c408 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Sat, 25 Nov 2017 12:55:27 +0100 Subject: [PATCH 09/38] TST: ignore cupy_tensors in doctests if cupy is not available --- odl/space/cupy_tensors.py | 42 +++++++++++++++++++++++--------------- odl/util/pytest_plugins.py | 5 +++++ 2 files changed, 31 insertions(+), 16 deletions(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index 69d5f5c4810..af81efe5a5b 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -431,6 +431,25 @@ def _lincomb_impl(a, x1, b, x2, out): lico(a, x1.data, b, x2.data, out.data) +def _python_scalar(arr): + """Convert a CuPy array to a Python scalar. + + The shape of the array must be ``()`` or ``(1,)``, otherwise an + error is raised. + """ + if arr.dtype.kind == 'f': + return float(cupy.asnumpy(arr)) + elif arr.dtype.kind == 'c': + return complex(cupy.asnumpy(arr)) + elif arr.dtype.kind in ('u', 'i'): + return int(cupy.asnumpy(arr)) + elif arr.dtype.kind == 'b': + return bool(cupy.asnumpy(arr)) + else: + raise RuntimeError("no conversion for dtype {}" + "".format(arr.dtype)) + + # --- Space and element classes --- # @@ -1337,15 +1356,7 @@ def __getitem__(self, indices): """ arr = self.data[indices] if arr.shape == (): - if arr.dtype.kind == 'f': - return float(cupy.asnumpy(arr)) - elif arr.dtype.kind == 'c': - return complex(cupy.asnumpy(arr)) - elif arr.dtype.kind in ('u', 'i'): - return int(cupy.asnumpy(arr)) - else: - raise RuntimeError("no conversion for dtype {}" - "".format(arr.dtype)) + return _python_scalar(arr) else: if is_numeric_dtype(self.dtype): weighting = self.space.weighting @@ -1898,7 +1909,7 @@ def eval_at_via_npy(*inputs, **kwargs): elif (ufunc in (np.add, np.multiply) and method in ('reduce', 'accumulate')): # These cases are implemented but not available as methods - # of cupy.ufunc. We do the manual assignment + # of cupy.ufunc. We map the implementation by hand. if ufunc == np.add: function = cupy.sum if method == 'reduce' else cupy.cumsum else: @@ -1906,10 +1917,10 @@ def eval_at_via_npy(*inputs, **kwargs): res = function(*inputs, **kwargs) - # Shortcut for scalar or no return value - if np.isscalar(res): + # Shortcut for scalar return value + if getattr(res, 'shape', ()) == (): # Happens for `reduce` with all axes - return res + return _python_scalar(res) # Wrap result if necessary (lazily) if out is None: @@ -2647,6 +2658,5 @@ def __init__(self, dist): if __name__ == '__main__': - if CUPY_AVAILABLE: - from odl.util.testutils import run_doctests - run_doctests() + from odl.util.testutils import run_doctests + run_doctests() diff --git a/odl/util/pytest_plugins.py b/odl/util/pytest_plugins.py index dbac0fef1a5..4054b21c71a 100644 --- a/odl/util/pytest_plugins.py +++ b/odl/util/pytest_plugins.py @@ -14,6 +14,7 @@ import os import odl +from odl.space.cupy_tensors import CUPY_AVAILABLE from odl.space.entry_points import tensor_space_impl_names from odl.trafos.backends import PYFFTW_AVAILABLE, PYWT_AVAILABLE from odl.util.testutils import simple_fixture @@ -80,6 +81,10 @@ def find_example_dirs(): collect_ignore.append( os.path.join(odl_root, 'odl', 'trafos', 'wavelet.py')) +if not CUPY_AVAILABLE: + collect_ignore.append( + os.path.join(odl_root, 'odl', 'space', 'cupy_tensors.py')) + # Remove duplicates collect_ignore = list(set(collect_ignore)) From dd7f6357dd5404b51b6513d20a0b85bebb4d2bf2 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Sat, 25 Nov 2017 15:55:59 +0100 Subject: [PATCH 10/38] MAINT: fix various numpy and cupy tensor things --- odl/space/cupy_tensors.py | 69 +++++++++++++++++++++++++++++---------- odl/space/npy_tensors.py | 2 +- 2 files changed, 52 insertions(+), 19 deletions(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index af81efe5a5b..5c9a749c9f4 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -50,6 +50,14 @@ operation='z = a * x + b * y;', name='lico') + all_equal = cupy.ReductionKernel(in_params='T x, T y', + out_params='R res', + map_expr='x == y', + reduce_expr='a && b', + post_map_expr='res = a', + identity='1', + name='all_equal') + def _fallback_scal(a, x): """Fallback implementation of ``scal`` when cuBLAS is not applicable.""" @@ -731,17 +739,21 @@ def element(self, inp=None, order=None): [ 4., 5., 6.]] ) """ - if order is None: - order_in = order - else: + if order is not None: + # Need to keep `order=None` intact order, order_in = str(order).upper(), order - - if order is not None and order not in ('C', 'F'): - raise ValueError("`order` {!r} not understood".format(order_in)) + if order not in ('C', 'F'): + raise ValueError('`order` {!r} not understood' + ''.format(order_in)) with cupy.cuda.Device(self.device): if inp is None: if order is None: + # TODO: remove when fix is available + # `order=None` is not understood by cupy 2.1.0, for new + # elements it means to use default order. See + # https://github.com/cupy/cupy/issues/590 + # (fixed on master) order = self.default_order arr = cupy.empty(self.shape, dtype=self.dtype, order=order) @@ -755,16 +767,22 @@ def element(self, inp=None, order=None): ''.format(self.shape, inp.shape)) if isinstance(inp, cupy.ndarray): - # Workaround for https://github.com/cupy/cupy/issues/590 - # TODO: remove when solved - if (inp.dtype == self.dtype and - inp.device.id == self.device): - arr = inp + # Copies to current device if necessary + # TODO: remove first case when `order=None` fix is + # available (see above) + if order is None: + arr = inp.astype(self.dtype, copy=False) else: - arr = inp.astype(self.dtype) + arr = inp.astype(self.dtype, order=order, copy=False) else: - arr = cupy.array(inp, copy=False, dtype=self.dtype, - ndmin=self.ndim, order=order) + # TODO: remove first case when `order=None` fix is + # available (see above) + if order is None: + arr = cupy.array(inp, copy=False, dtype=self.dtype, + ndmin=self.ndim) + else: + arr = cupy.array(inp, copy=False, dtype=self.dtype, + ndmin=self.ndim, order=order) # If the result has a 0 stride, make a copy since it would # produce all kinds of nasty problems. This happens for e.g. @@ -1206,6 +1224,9 @@ def asarray(self, out=None): if out is None: return cupy.asnumpy(self.data) else: + if not isinstance(out, np.ndarray): + raise TypeError('`out` must be a `numpy.ndarray´, got type ' + '{}'.format(type(out))) if out.shape != self.shape: raise ValueError('`out` must have shape {}, got shape {}' ''.format(self.shape, out.shape)) @@ -1219,7 +1240,7 @@ def asarray(self, out=None): @property def data_ptr(self): - """A raw pointer to the data container. + """Memory address of the data container as 64-bit integer. Examples -------- @@ -1272,7 +1293,8 @@ def __eq__(self, other): elif other not in self.space: return False else: - return bool((self.data == other.data).all()) + out = cupy.empty((), dtype=bool) + return bool(all_equal(self.data, other.data, out)) def copy(self): """Create an identical (deep) copy of this tensor. @@ -1321,8 +1343,8 @@ def __getitem__(self, indices): >>> x[::2] rn(3, impl='cupy').element([ 1., 3., 5.]) - The returned views are writable, so modificatons alter the - original array: + If integers and slices are used, the returned "views" are writable, + hence modificatons alter the original array: >>> view = x[1:4] >>> view[:] = -1 @@ -1331,6 +1353,17 @@ def __getitem__(self, indices): >>> x rn(5, impl='cupy').element([ 1., -1., -1., -1., 5.]) + For lists, index arrays or boolean arrays, a copy is made, i.e., + changes do not affect the original array: + + >>> idcs = [0, 3, 2, 1, 2] + >>> x_part = x[[0, 3, 2, 1, 2]] + >>> x_part + rn(5, impl='cupy').element([ 1., 4., 3., 2., 3.]) + >>> x_part[:] = 0 + >>> x + rn(5, impl='cupy').element([ 1., 2., 3., 4., 5.]) + Multi-indexing is also directly supported: >>> tensors = odl.rn((2, 3), impl='cupy') diff --git a/odl/space/npy_tensors.py b/odl/space/npy_tensors.py index f0e04bc7795..ae75bf36ab7 100644 --- a/odl/space/npy_tensors.py +++ b/odl/space/npy_tensors.py @@ -927,7 +927,7 @@ def astype(self, dtype): @property def data_ptr(self): - """A raw pointer to the data container of ``self``. + """Memory address of the data container as 64-bit integer. Examples -------- From 5ab4756e09c425a3a10ea2f9b2ddcc24a68d9cfe Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Sat, 25 Nov 2017 15:56:17 +0100 Subject: [PATCH 11/38] ENH: allow indexing of CupyTensor with CupyTensor --- odl/space/cupy_tensors.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index 5c9a749c9f4..c1d1d8b2b4e 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -1387,6 +1387,9 @@ def __getitem__(self, indices): [ 0., 5., 0.]] ) """ + if isinstance(indices, CupyTensor): + indices = indices.data + arr = self.data[indices] if arr.shape == (): return _python_scalar(arr) From fefb72289645e330605bdbabdf5ea960ba71aa3a Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Sat, 25 Nov 2017 21:06:04 +0100 Subject: [PATCH 12/38] MAINT: minor stuff --- odl/space/cupy_tensors.py | 5 ++--- odl/space/npy_tensors.py | 5 +++-- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index c1d1d8b2b4e..b652ba0af87 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -1141,8 +1141,7 @@ def default_dtype(field=None): - ``RealNumbers()`` or ``None`` : ``np.dtype('float64')`` - ``ComplexNumbers()`` : ``np.dtype('complex128')`` - These choices correspond to the defaults of the ``cupy`` - library. + These choices correspond to the defaults of the CuPy library. """ if field is None or field == RealNumbers(): return np.dtype('float64') @@ -1225,7 +1224,7 @@ def asarray(self, out=None): return cupy.asnumpy(self.data) else: if not isinstance(out, np.ndarray): - raise TypeError('`out` must be a `numpy.ndarray´, got type ' + raise TypeError('`out` must be a `numpy.ndarray`, got type ' '{}'.format(type(out))) if out.shape != self.shape: raise ValueError('`out` must have shape {}, got shape {}' diff --git a/odl/space/npy_tensors.py b/odl/space/npy_tensors.py index ae75bf36ab7..219b3f8f8cf 100644 --- a/odl/space/npy_tensors.py +++ b/odl/space/npy_tensors.py @@ -502,9 +502,10 @@ def default_dtype(field=None): dtype : `numpy.dtype` Numpy data type specifier. The returned defaults are: - ``RealNumbers()`` : ``np.dtype('float64')`` + - ``RealNumbers()`` : ``np.dtype('float64')`` + - ``ComplexNumbers()`` : ``np.dtype('complex128')`` - ``ComplexNumbers()`` : ``np.dtype('complex128')`` + These choices correspond to the defaults of the NumPy library. """ if field is None or field == RealNumbers(): return np.dtype('float64') From 6648425d6e765c9234caadbe3e7e02326aa9c37f Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Sat, 25 Nov 2017 23:36:44 +0100 Subject: [PATCH 13/38] MAINT: change weighted inf-norm to ignore weight --- odl/space/npy_tensors.py | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/odl/space/npy_tensors.py b/odl/space/npy_tensors.py index 219b3f8f8cf..f3a87751bc5 100644 --- a/odl/space/npy_tensors.py +++ b/odl/space/npy_tensors.py @@ -2021,7 +2021,6 @@ def _pnorm_diagweight(x, p, w): # BLAS dot or nrm2 xp = np.abs(x.data.ravel(order)) if p == float('inf'): - xp *= w.ravel(order) return np.max(xp) else: xp = np.power(xp, p, out=xp) @@ -2276,7 +2275,7 @@ def norm(self, x): if self.exponent == 2.0: return float(np.sqrt(self.const) * _norm_default(x)) elif self.exponent == float('inf'): - return float(self.const * _pnorm_default(x, self.exponent)) + return float(_pnorm_default(x, self.exponent)) else: return float((self.const ** (1 / self.exponent) * _pnorm_default(x, self.exponent))) @@ -2297,7 +2296,7 @@ def dist(self, x1, x2): if self.exponent == 2.0: return float(np.sqrt(self.const) * _norm_default(x1 - x2)) elif self.exponent == float('inf'): - return float(self.const * _pnorm_default(x1 - x2, self.exponent)) + return float(_pnorm_default(x1 - x2, self.exponent)) else: return float((self.const ** (1 / self.exponent) * _pnorm_default(x1 - x2, self.exponent))) From a1cd4829d85e0232afacddd4c7630ba8a7d2f062 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Sat, 25 Nov 2017 23:38:09 +0100 Subject: [PATCH 14/38] ENH: make comparison of cupy arrays in tests faster --- odl/util/testutils.py | 52 +++++++++++++++++++++++++++++++++++-------- 1 file changed, 43 insertions(+), 9 deletions(-) diff --git a/odl/util/testutils.py b/odl/util/testutils.py index 9339b65e081..4efa4f6be74 100644 --- a/odl/util/testutils.py +++ b/odl/util/testutils.py @@ -91,6 +91,19 @@ def dtype_tol(dtype, default=None): def all_equal(iter1, iter2): """Return ``True`` if all elements in ``a`` and ``b`` are equal.""" + # Transfer cupy arrays to CPU for faster comparison + from odl.space.cupy_tensors import CUPY_AVAILABLE, CupyTensor, cupy + if CUPY_AVAILABLE: + if isinstance(iter1, CupyTensor): + iter1 = iter1.asarray() + elif isinstance(iter1, cupy.ndarray): + iter1 = cupy.asnumpy(iter1) + + if isinstance(iter2, CupyTensor): + iter2 = iter2.asarray() + elif isinstance(iter2, cupy.ndarray): + iter2 = cupy.asnumpy(iter2) + # Direct comparison for scalars, tuples or lists try: if iter1 == iter2: @@ -137,6 +150,19 @@ def all_almost_equal_array(v1, v2, ndigits): def all_almost_equal(iter1, iter2, ndigits=None): """Return ``True`` if all elements in ``a`` and ``b`` are almost equal.""" + # Transfer cupy arrays to CPU for faster comparison + from odl.space.cupy_tensors import CUPY_AVAILABLE, CupyTensor, cupy + if CUPY_AVAILABLE: + if isinstance(iter1, CupyTensor): + iter1 = iter1.asarray() + elif isinstance(iter1, cupy.ndarray): + iter1 = cupy.asnumpy(iter1) + + if isinstance(iter2, CupyTensor): + iter2 = iter2.asarray() + elif isinstance(iter2, cupy.ndarray): + iter2 = cupy.asnumpy(iter2) + try: if iter1 is iter2 or iter1 == iter2: return True @@ -323,9 +349,11 @@ def noise_array(space): odl.set.space.LinearSpace.examples : Examples of elements typical to the space. """ - from odl.space import ProductSpace + from odl.space.pspace import ProductSpace + from odl.space.cupy_tensors import cupy + if isinstance(space, ProductSpace): - return np.array([noise_array(si) for si in space]) + return [noise_array(si) for si in space] else: if space.dtype == bool: # TODO(kohr-h): use `randint(..., dtype=bool)` from Numpy 1.11 on @@ -342,7 +370,13 @@ def noise_array(space): else: raise ValueError('bad dtype {}'.format(space.dtype)) - return arr.astype(space.dtype, copy=False) + arr = arr.astype(space.dtype, copy=False) + if space.impl == 'numpy': + return arr + elif space.impl == 'cupy': + return cupy.asarray(arr) + else: + raise RuntimeError('bad `impl` {!r}'.format(space.impl)) def noise_element(space): @@ -557,23 +591,23 @@ class ProgressBar(object): Usage: >>> progress = ProgressBar('Reading data', 10) - \rReading data: [ ] Starting + Reading data: [ ] Starting >>> progress.update(4) #halfway, zero indexing - \rReading data: [############### ] 50.0% + Reading data: [############### ] 50.0% Multi-indices, from slowest to fastest: >>> progress = ProgressBar('Reading data', 10, 10) - \rReading data: [ ] Starting + Reading data: [ ] Starting >>> progress.update(9, 8) - \rReading data: [############################# ] 99.0% + Reading data: [############################# ] 99.0% Supports simply calling update, which moves the counter forward: >>> progress = ProgressBar('Reading data', 10, 10) - \rReading data: [ ] Starting + Reading data: [ ] Starting >>> progress.update() - \rReading data: [ ] 1.0% + Reading data: [ ] 1.0% """ def __init__(self, text='progress', *njobs): From cad427c4cecdcbde655fb445972c7400b1e12eb1 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Sat, 25 Nov 2017 23:39:01 +0100 Subject: [PATCH 15/38] WIP: fix tensor tests for cupy --- odl/space/cupy_tensors.py | 2 + odl/test/space/tensors_test.py | 245 ++++++++++++++++++--------------- 2 files changed, 139 insertions(+), 108 deletions(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index b652ba0af87..3a3615c6a3e 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -28,6 +28,7 @@ try: import cupy except ImportError: + cupy = None CUPY_AVAILABLE = False else: _maj = int(cupy.__version__.split('.')[0]) @@ -1356,6 +1357,7 @@ def __getitem__(self, indices): changes do not affect the original array: >>> idcs = [0, 3, 2, 1, 2] + >>> x = r5.element([1, 2, 3, 4, 5]) >>> x_part = x[[0, 3, 2, 1, 2]] >>> x_part rn(5, impl='cupy').element([ 1., 4., 3., 2., 3.]) diff --git a/odl/test/space/tensors_test.py b/odl/test/space/tensors_test.py index 24bc061ba20..8f8d92b886e 100644 --- a/odl/test/space/tensors_test.py +++ b/odl/test/space/tensors_test.py @@ -27,15 +27,12 @@ CupyTensorSpaceConstWeighting, CupyTensorSpaceArrayWeighting, CupyTensorSpaceCustomInner, CupyTensorSpaceCustomNorm, CupyTensorSpaceCustomDist, - CUPY_AVAILABLE) + CUPY_AVAILABLE, cupy) from odl.util.testutils import ( all_almost_equal, all_equal, simple_fixture, noise_array, noise_element, noise_elements) from odl.util.ufuncs import UFUNCS -if CUPY_AVAILABLE: - import cupy - # --- Test helpers --- # @@ -44,28 +41,27 @@ parse_version('1.13')) -# Functions to return arrays and classes corresponding to impls. Extend +# Functions to return arrays, classes etc. corresponding to impls. Extend # when a new impl is available. -def _pos_array(space): - """Create a Numpy array with positive real entries for ``space``.""" - arr = np.abs(noise_array(space)) + 0.1 - if space.impl == 'numpy': - return arr - elif space.impl == 'cupy': - return cupy.asarray(arr) +def _module(impl): + """Return the array module for ``impl``.""" + if impl == 'numpy': + return np + elif impl == 'cupy': + return cupy else: assert False +def _pos_array(space): + """Create an array with positive real entries for ``space``.""" + return _module(space.impl).abs(noise_array(space)) + 0.1 + + def _array_cls(impl): """Return the array class for given impl.""" - if impl == 'numpy': - return np.ndarray - elif impl == 'cupy': - return cupy.ndarray - else: - assert False + return _module(impl).ndarray def _odl_tensor_cls(impl): @@ -135,10 +131,14 @@ def weight(request): @pytest.fixture(scope='module') -def tspace(odl_floating_dtype, odl_tspace_impl): - impl = odl_tspace_impl - dtype = odl_floating_dtype - return odl.tensor_space(shape=(3, 4), dtype=dtype, impl=impl) +def tspace(floating_dtype, odl_tspace_impl): + cls = odl.space.entry_points.tensor_space_impl(odl_tspace_impl) + if floating_dtype not in cls.available_dtypes(): + pytest.skip('dtype {} not supported by impl {!r}' + ''.format(floating_dtype, odl_tspace_impl)) + else: + return odl.tensor_space(shape=(3, 4), dtype=floating_dtype, + impl=odl_tspace_impl) # --- Space classes --- # @@ -254,16 +254,17 @@ def test_init_cupy_tspace(): odl.rn((3, 4), weighting=weight_arr, impl='cupy') -def test_init_tspace_weighting(weight, exponent, tspace_impl): +def test_init_tspace_weighting(weight, exponent, odl_tspace_impl): """Test if weightings during init give the correct weighting classes.""" - if tspace_impl == 'cupy' and isinstance(weight, np.ndarray): + impl = odl_tspace_impl + if impl == 'cupy' and isinstance(weight, np.ndarray): # Need cast before using in space creation since # ArrayWeighting.__eq__ uses `arr1 is arr2` to check arrays weight = cupy.asarray(weight) - elif isinstance(weight, _array_cls(tspace_impl)): - weighting_cls = _weighting_cls(tspace_impl, 'array') + elif isinstance(weight, _array_cls(impl)): + weighting_cls = _weighting_cls(impl, 'array') else: - weighting_cls = _weighting_cls(tspace_impl, 'const') + weighting_cls = _weighting_cls(impl, 'const') weighting = weighting_cls(weight, exponent) @@ -302,6 +303,9 @@ def test_properties(odl_tspace_impl): assert x.itemsize == 4 assert x.nbytes == 4 * 3 * 4 + if impl == 'cupy': + assert x.device == space.device == cupy.cuda.get_device_id() + def test_size(odl_tspace_impl): """Test that size handles corner cases appropriately.""" @@ -367,14 +371,15 @@ def test_element(tspace, odl_elem_order): assert elem.data.flags[order + '_CONTIGUOUS'] # From pointer - arr_c_ptr = arr_c.ctypes.data - elem = tspace.element(data_ptr=arr_c_ptr, order='C') - assert all_equal(elem, arr_c) - assert np.may_share_memory(elem.data, arr_c) - arr_f_ptr = arr_f.ctypes.data - elem = tspace.element(data_ptr=arr_f_ptr, order='F') - assert all_equal(elem, arr_f) - assert np.may_share_memory(elem.data, arr_f) + if tspace.impl == 'numpy': + arr_c_ptr = arr_c.ctypes.data + elem = tspace.element(data_ptr=arr_c_ptr, order='C') + assert all_equal(elem, arr_c) + assert np.may_share_memory(elem.data, arr_c) + arr_f_ptr = arr_f.ctypes.data + elem = tspace.element(data_ptr=arr_f_ptr, order='F') + assert all_equal(elem, arr_f) + assert np.may_share_memory(elem.data, arr_f) # Check errors with pytest.raises(ValueError): @@ -746,7 +751,7 @@ def test_norm(tspace): """Test the norm method against numpy.linalg.norm.""" xarr, x = noise_elements(tspace) - correct_norm = np.linalg.norm(xarr.ravel()) + correct_norm = _module(tspace.impl).linalg.norm(xarr.ravel()) assert tspace.norm(x) == pytest.approx(correct_norm) assert x.norm() == pytest.approx(correct_norm) @@ -760,12 +765,18 @@ def test_norm_exceptions(tspace): tspace.norm(other_x) -def test_pnorm(exponent): +def test_pnorm(exponent, odl_tspace_impl): """Test the norm method with p!=2 against numpy.linalg.norm.""" - for tspace in (odl.rn((3, 4), exponent=exponent), - odl.cn((3, 4), exponent=exponent)): + impl = odl_tspace_impl + spaces = [odl.rn((3, 4), exponent=exponent, impl=impl)] + cls = odl.space.entry_points.tensor_space_impl(impl) + if complex in cls.available_dtypes(): + spaces.append(odl.cn((3, 4), exponent=exponent, impl=impl)) + + for tspace in spaces: xarr, x = noise_elements(tspace) - correct_norm = np.linalg.norm(xarr.ravel(), ord=exponent) + correct_norm = _module(impl).linalg.norm( + xarr.ravel(), ord=exponent) assert tspace.norm(x) == pytest.approx(correct_norm) assert x.norm() == pytest.approx(correct_norm) @@ -775,7 +786,7 @@ def test_dist(tspace): """Test the dist method against numpy.linalg.norm of the difference.""" [xarr, yarr], [x, y] = noise_elements(tspace, n=2) - correct_dist = np.linalg.norm((xarr - yarr).ravel()) + correct_dist = _module(tspace.impl).linalg.norm((xarr - yarr).ravel()) assert tspace.dist(x, y) == pytest.approx(correct_dist) assert x.dist(y) == pytest.approx(correct_dist) @@ -800,10 +811,12 @@ def test_pdist(odl_tspace_impl, exponent): cls = odl.space.entry_points.tensor_space_impl(impl) if complex in cls.available_dtypes(): spaces.append(odl.cn((3, 4), exponent=exponent, impl=impl)) + for space in spaces: [xarr, yarr], [x, y] = noise_elements(space, n=2) - correct_dist = np.linalg.norm((xarr - yarr).ravel(), ord=exponent) + correct_dist = _module(impl).linalg.norm( + (xarr - yarr).ravel(), ord=exponent) assert space.dist(x, y) == pytest.approx(correct_dist) assert x.dist(y) == pytest.approx(correct_dist) @@ -830,13 +843,12 @@ def test_element_getitem(odl_tspace_impl, getitem_indices): assert sliced_spc.shape == sliced_shape assert sliced_spc.dtype == space.dtype assert sliced_spc.exponent == space.exponent - assert sliced_spc.weighting == space.weighting # Check that we have a view that manipulates the original array # (or not, depending on indexing style) x_arr_sliced[:] = 0 x_sliced[:] = 0 - assert all_equal(x_arr, x) + assert all_equal(x, x_arr) def test_element_setitem(odl_tspace_impl, setitem_indices): @@ -873,19 +885,19 @@ def test_element_getitem_bool_array(odl_tspace_impl): space = odl.tensor_space((2, 3, 4), dtype='float32', exponent=1, weighting=2, impl=impl) bool_space = odl.tensor_space((2, 3, 4), dtype=bool) + x_arr, x = noise_elements(space) cond_arr, cond = noise_elements(bool_space) x_arr_sliced = x_arr[cond_arr] x_sliced = x[cond] - assert all_equal(x_arr_sliced, x_sliced) + assert all_equal(x_arr_sliced.asarray(), x_sliced.asarray()) # Check that the space properties are preserved sliced_spc = x_sliced.space assert sliced_spc.shape == x_arr_sliced.shape assert sliced_spc.dtype == space.dtype assert sliced_spc.exponent == space.exponent - assert sliced_spc.weighting == space.weighting def test_element_setitem_bool_array(odl_tspace_impl): @@ -1130,7 +1142,7 @@ def test_array_weighting_array_is_valid(odl_tspace_impl): # Invalid weight_arr[0] = 0 - weighting_arr = NumpyTensorSpaceArrayWeighting(weight_arr) + weighting_arr = weighting_cls(weight_arr) assert not weighting_arr.is_valid() @@ -1201,54 +1213,65 @@ def test_array_weighting_inner(tspace): [xarr, yarr], [x, y] = noise_elements(tspace, 2) weight_arr = _pos_array(tspace) - weighting = NumpyTensorSpaceArrayWeighting(weight_arr) + weighting_cls = _weighting_cls(tspace.impl, 'array') + weighting = weighting_cls(weight_arr) - true_inner = np.vdot(yarr, xarr * weight_arr) - assert weighting.inner(x, y) == pytest.approx(true_inner) + if tspace.impl == 'numpy': + true_inner = np.vdot(yarr, xarr * weight_arr) + elif tspace.impl == 'cupy': + true_inner = cupy.vdot(yarr, xarr * weight_arr) + else: + assert False + + assert weighting.inner(x, y) == pytest.approx(true_inner, rel=1e-2) # Exponent != 2 -> no inner product, should raise with pytest.raises(NotImplementedError): - NumpyTensorSpaceArrayWeighting(weight_arr, exponent=1.0).inner(x, y) + weighting_cls(weight_arr, exponent=1.0).inner(x, y) def test_array_weighting_norm(tspace, exponent): """Test norm in a weighted space.""" - rtol = np.sqrt(np.finfo(tspace.dtype).resolution) xarr, x = noise_elements(tspace) weight_arr = _pos_array(tspace) - weighting = NumpyTensorSpaceArrayWeighting(weight_arr, exponent=exponent) + weighting_cls = _weighting_cls(tspace.impl, 'array') + weighting = weighting_cls(weight_arr, exponent=exponent) + + if tspace.impl == 'numpy': + norm = np.linalg.norm + elif tspace.impl == 'cupy': + norm = cupy.linalg.norm + else: + assert False if exponent == float('inf'): - true_norm = np.linalg.norm( - (weight_arr * xarr).ravel(), - ord=float('inf')) + true_norm = float(norm((xarr).ravel(), ord=float('inf'))) else: - true_norm = np.linalg.norm( - (weight_arr ** (1 / exponent) * xarr).ravel(), - ord=exponent) + true_norm = float(norm((weight_arr ** (1 / exponent) * xarr).ravel(), + ord=exponent)) - assert weighting.norm(x) == pytest.approx(true_norm, rel=rtol) + assert weighting.norm(x) == pytest.approx(true_norm, rel=1e-2) def test_array_weighting_dist(tspace, exponent): """Test dist product in a weighted space.""" - rtol = np.sqrt(np.finfo(tspace.dtype).resolution) [xarr, yarr], [x, y] = noise_elements(tspace, n=2) weight_arr = _pos_array(tspace) - weighting = NumpyTensorSpaceArrayWeighting(weight_arr, exponent=exponent) + weighting_cls = _weighting_cls(tspace.impl, 'array') + weighting = weighting_cls(weight_arr, exponent=exponent) if exponent == float('inf'): - true_dist = np.linalg.norm( - (weight_arr * (xarr - yarr)).ravel(), + true_dist = _module(tspace.impl).linalg.norm( + ((xarr - yarr)).ravel(), ord=float('inf')) else: - true_dist = np.linalg.norm( + true_dist = _module(tspace.impl).linalg.norm( (weight_arr ** (1 / exponent) * (xarr - yarr)).ravel(), ord=exponent) - assert weighting.dist(x, y) == pytest.approx(true_dist, rel=rtol) + assert weighting.dist(x, y) == pytest.approx(true_dist, rel=1e-2) def test_const_weighting_init(odl_tspace_impl, exponent): @@ -1312,13 +1335,14 @@ def test_const_weighting_inner(tspace): [xarr, yarr], [x, y] = noise_elements(tspace, 2) constant = 1.5 - true_result_const = constant * np.vdot(yarr, xarr) + true_result_const = constant * _module(tspace.impl).vdot(yarr, xarr) - w_const = NumpyTensorSpaceConstWeighting(constant) - assert w_const.inner(x, y) == pytest.approx(true_result_const) + weighting_cls = _weighting_cls(tspace.impl, 'const') + w_const = weighting_cls(constant) + assert w_const.inner(x, y) == pytest.approx(true_result_const, rel=1e-2) # Exponent != 2 -> no inner - w_const = NumpyTensorSpaceConstWeighting(constant, exponent=1) + w_const = weighting_cls(constant, exponent=1) with pytest.raises(NotImplementedError): w_const.inner(x, y) @@ -1329,13 +1353,15 @@ def test_const_weighting_norm(tspace, exponent): constant = 1.5 if exponent == float('inf'): - factor = constant + factor = 1.0 else: factor = constant ** (1 / exponent) - true_norm = factor * np.linalg.norm(xarr.ravel(), ord=exponent) + true_norm = factor * _module(tspace.impl).linalg.norm( + xarr.ravel(), ord=exponent) - w_const = NumpyTensorSpaceConstWeighting(constant, exponent=exponent) - assert w_const.norm(x) == pytest.approx(true_norm) + weighting_cls = _weighting_cls(tspace.impl, 'const') + w_const = weighting_cls(constant, exponent=exponent) + assert w_const.norm(x) == pytest.approx(true_norm, rel=1e-2) def test_const_weighting_dist(tspace, exponent): @@ -1344,57 +1370,59 @@ def test_const_weighting_dist(tspace, exponent): constant = 1.5 if exponent == float('inf'): - factor = constant + factor = 1.0 else: factor = constant ** (1 / exponent) - true_dist = factor * np.linalg.norm((xarr - yarr).ravel(), ord=exponent) + true_dist = factor * _module(tspace.impl).linalg.norm( + (xarr - yarr).ravel(), ord=exponent) - w_const = NumpyTensorSpaceConstWeighting(constant, exponent=exponent) - assert w_const.dist(x, y) == pytest.approx(true_dist) + weighting_cls = _weighting_cls(tspace.impl, 'const') + w_const = weighting_cls(constant, exponent=exponent) + assert w_const.dist(x, y) == pytest.approx(true_dist, rel=1e-2) def test_custom_inner(tspace): """Test weighting with a custom inner product.""" - rtol = np.sqrt(np.finfo(tspace.dtype).resolution) - [xarr, yarr], [x, y] = noise_elements(tspace, 2) def inner(x, y): - return np.vdot(y, x) + return _module(tspace.impl).vdot(y, x) - w = NumpyTensorSpaceCustomInner(inner) - w_same = NumpyTensorSpaceCustomInner(inner) - w_other = NumpyTensorSpaceCustomInner(np.dot) + weighting_cls = _weighting_cls(tspace.impl, 'inner') + w = weighting_cls(inner) + w_same = weighting_cls(inner) + w_other = weighting_cls(_module(tspace.impl).dot) assert w == w assert w == w_same assert w != w_other true_inner = inner(xarr, yarr) - assert w.inner(x, y) == pytest.approx(true_inner) + assert w.inner(x, y) == pytest.approx(true_inner, rel=1e-2) - true_norm = np.linalg.norm(xarr.ravel()) - assert w.norm(x) == pytest.approx(true_norm) + true_norm = _module(tspace.impl).linalg.norm(xarr.ravel()) + assert w.norm(x) == pytest.approx(true_norm, rel=1e-2) - true_dist = np.linalg.norm((xarr - yarr).ravel()) - assert w.dist(x, y) == pytest.approx(true_dist, rel=rtol) + true_dist = _module(tspace.impl).linalg.norm((xarr - yarr).ravel()) + assert w.dist(x, y) == pytest.approx(true_dist, rel=1e-2) with pytest.raises(TypeError): - NumpyTensorSpaceCustomInner(1) + weighting_cls(1) def test_custom_norm(tspace): """Test weighting with a custom norm.""" [xarr, yarr], [x, y] = noise_elements(tspace, 2) - norm = np.linalg.norm + norm = _module(tspace.impl).linalg.norm def other_norm(x): - return np.linalg.norm(x, ord=1) + return _module(tspace.impl).linalg.norm(x, ord=1) - w = NumpyTensorSpaceCustomNorm(norm) - w_same = NumpyTensorSpaceCustomNorm(norm) - w_other = NumpyTensorSpaceCustomNorm(other_norm) + weighting_cls = _weighting_cls(tspace.impl, 'norm') + w = weighting_cls(norm) + w_same = weighting_cls(norm) + w_other = weighting_cls(other_norm) assert w == w assert w == w_same @@ -1403,14 +1431,14 @@ def other_norm(x): with pytest.raises(NotImplementedError): w.inner(x, y) - true_norm = np.linalg.norm(xarr.ravel()) - assert w.norm(x) == pytest.approx(true_norm) + true_norm = _module(tspace.impl).linalg.norm(xarr.ravel()) + assert w.norm(x) == pytest.approx(true_norm, rel=1e-2) - true_dist = np.linalg.norm((xarr - yarr).ravel()) - assert w.dist(x, y) == pytest.approx(true_dist) + true_dist = _module(tspace.impl).linalg.norm((xarr - yarr).ravel()) + assert w.dist(x, y) == pytest.approx(true_dist, rel=1e-2) with pytest.raises(TypeError): - NumpyTensorSpaceCustomNorm(1) + weighting_cls(1) def test_custom_dist(tspace): @@ -1418,14 +1446,15 @@ def test_custom_dist(tspace): [xarr, yarr], [x, y] = noise_elements(tspace, 2) def dist(x, y): - return np.linalg.norm(x - y) + return _module(tspace.impl).linalg.norm(x - y) def other_dist(x, y): - return np.linalg.norm(x - y, ord=1) + return _module(tspace.impl).linalg.norm(x - y, ord=1) - w = NumpyTensorSpaceCustomDist(dist) - w_same = NumpyTensorSpaceCustomDist(dist) - w_other = NumpyTensorSpaceCustomDist(other_dist) + weighting_cls = _weighting_cls(tspace.impl, 'dist') + w = weighting_cls(dist) + w_same = weighting_cls(dist) + w_other = weighting_cls(other_dist) assert w == w assert w == w_same @@ -1437,11 +1466,11 @@ def other_dist(x, y): with pytest.raises(NotImplementedError): w.norm(x) - true_dist = np.linalg.norm((xarr - yarr).ravel()) - assert w.dist(x, y) == pytest.approx(true_dist) + true_dist = _module(tspace.impl).linalg.norm((xarr - yarr).ravel()) + assert w.dist(x, y) == pytest.approx(true_dist, rel=1e-2) with pytest.raises(TypeError): - NumpyTensorSpaceCustomDist(1) + weighting_cls(1) # --- Ufuncs & Reductions --- # From 658051f6625665481711112ac724dfc8baf56856 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Sun, 26 Nov 2017 17:05:04 +0100 Subject: [PATCH 16/38] MAINT: adapt all_almost_equal to cupy --- odl/util/testutils.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/odl/util/testutils.py b/odl/util/testutils.py index 4efa4f6be74..deb4584ecf4 100644 --- a/odl/util/testutils.py +++ b/odl/util/testutils.py @@ -166,7 +166,7 @@ def all_almost_equal(iter1, iter2, ndigits=None): try: if iter1 is iter2 or iter1 == iter2: return True - except ValueError: + except (ValueError, TypeError): pass if iter1 is None and iter2 is None: From 8a7dc33fd8cf7a4ee9c6058ba1483500ee38c655 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Sun, 26 Nov 2017 17:06:00 +0100 Subject: [PATCH 17/38] WIP: fix cupy tensor tests --- odl/space/cupy_tensors.py | 42 ++++++--- odl/test/space/tensors_test.py | 162 ++++++++++++++++++--------------- 2 files changed, 119 insertions(+), 85 deletions(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index 3a3615c6a3e..1195ac726c7 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -22,8 +22,8 @@ Weighting, ArrayWeighting, ConstWeighting, CustomInner, CustomNorm, CustomDist) from odl.util import ( - array_str, dtype_str, is_floating_dtype, is_numeric_dtype, real_dtype, - signature_string, indent) + array_str, dtype_str, is_floating_dtype, is_numeric_dtype, is_real_dtype, + real_dtype, signature_string, indent) try: import cupy @@ -2163,7 +2163,7 @@ def __ipow__(self, other): except TypeError: pass - self.ufuncs.power(self.data, other, out=self.data) + self.ufuncs.power(other, out=self) return self @@ -2216,7 +2216,7 @@ def _weighting(weights, exponent): reduce_expr='a + b', post_map_expr='res = a', identity='0', - name='nrm1w') + name='nrm1') nrm1w = cupy.ReductionKernel(in_params='T x, W w', out_params='R res', @@ -2228,7 +2228,7 @@ def _weighting(weights, exponent): nrm2 = cupy.ReductionKernel(in_params='T x', out_params='R res', - map_expr='x * x', + map_expr='abs(x * x)', reduce_expr='a + b', post_map_expr='res = sqrt(a)', identity='0', @@ -2236,7 +2236,7 @@ def _weighting(weights, exponent): nrm2w = cupy.ReductionKernel(in_params='T x, W w', out_params='R res', - map_expr='x * x * w', + map_expr='abs(x * x) * w', reduce_expr='a + b', post_map_expr='res = sqrt(a)', identity='0', @@ -2291,7 +2291,7 @@ def _weighting(weights, exponent): name='dist1') dist1w = cupy.ReductionKernel(in_params='T x, T y, W w', - out_params='T res', + out_params='R res', map_expr='abs(x - y) * w', reduce_expr='a + b', post_map_expr='res = a', @@ -2360,7 +2360,8 @@ def __init__(self, array, exponent=2.0): Parameters ---------- array : `array-like`, one-dim. - Weighting array of the inner product, norm and distance. + Weighting array of the inner product, norm and distance, + must have real data type and should only have positive entries. exponent : positive float Exponent of the norm. For values other than 2.0, the inner product is not defined. @@ -2369,6 +2370,11 @@ def __init__(self, array, exponent=2.0): array = array.data elif not isinstance(array, cupy.ndarray): array = cupy.array(array, copy=False) + + if not is_real_dtype(array.dtype): + raise ValueError('`array.dtype` must be real, got {}' + ''.format(dtype_str(array.dtype))) + super(CupyTensorSpaceArrayWeighting, self).__init__( array, impl='cupy', exponent=exponent) @@ -2390,7 +2396,10 @@ def inner(self, x1, x2): 'exponent != 2 (got {})' ''.format(self.exponent)) else: - return x1.space.field.element(dotw(x1.data, x2.data, self.array)) + # complex(cupy_complex_scalar) not implemented, see + # https://github.com/cupy/cupy/issues/782 + return x1.space.field.element( + cupy.asnumpy(dotw(x2.data.conj(), x1.data, self.array))) def norm(self, x): """Return the weighted norm of a tensor. @@ -2526,8 +2535,14 @@ def inner(self, x1, x2): 'exponent != 2 (got {})' ''.format(self.exponent)) else: - return x1.space.field.element( - self.const * cupy.dot(x1.data, x2.data)) + if np.issubsctype(x1.dtype, np.complexfloating): + dot = cupy.vdot(x2.data, x1.data) + else: + dot = cupy.dot(x2.data, x1.data) + + # complex(cupy_complex_scalar) not implemented, see + # https://github.com/cupy/cupy/issues/782 + return x1.space.field.element(self.const * cupy.asnumpy(dot)) def norm(self, x): """Return the constant-weighted norm of a tensor. @@ -2570,7 +2585,8 @@ def norm(self, x): pass else: norm = _cublas_nrm2( - x.data.device.cublas_handle, x.size, x.data_ptr, incx) + x.data.data.device.cublas_handle, x.size, + x.data.data.ptr, incx) return float(np.sqrt(self.const) * norm) # Cannot use cuBLAS, fall back to custom kernel @@ -2581,7 +2597,7 @@ def norm(self, x): return float(nrmneginf(x.data, out)) else: return float(self.const ** (1 / self.exponent) * - nrmp(x, self.exponent, out)) + nrmp(x.data, self.exponent, out)) def dist(self, x1, x2): """Return the weighted distance between two tensors. diff --git a/odl/test/space/tensors_test.py b/odl/test/space/tensors_test.py index 8f8d92b886e..b884e5cc0f0 100644 --- a/odl/test/space/tensors_test.py +++ b/odl/test/space/tensors_test.py @@ -54,16 +54,36 @@ def _module(impl): assert False -def _pos_array(space): - """Create an array with positive real entries for ``space``.""" - return _module(space.impl).abs(noise_array(space)) + 0.1 - - def _array_cls(impl): """Return the array class for given impl.""" return _module(impl).ndarray +def _as_numpy(array): + """Return a numpy.ndarray from the given array.""" + if isinstance(array, np.ndarray): + return array + elif isinstance(array, cupy.ndarray): + return cupy.asnumpy(array) + else: + assert False + + +def _data_ptr(array): + """Return the memory address of the given array (depending on impl).""" + if isinstance(array, np.ndarray): + return array.ctypes.data + elif isinstance(array, cupy.ndarray): + return array.data.ptr + else: + assert False + + +def _pos_array(space): + """Create an array with positive real entries for ``space``.""" + return _module(space.impl).abs(noise_array(space)) + 0.1 + + def _odl_tensor_cls(impl): """Return the ODL tensor class for given impl.""" if impl == 'numpy': @@ -344,8 +364,8 @@ def test_element(tspace, odl_elem_order): else: assert elem.data.flags[order + '_CONTIGUOUS'] - # From Numpy array (C order) - arr_c = np.random.rand(*tspace.shape).astype(tspace.dtype) + # From array (C order) + arr_c = _module(tspace.impl).ascontiguousarray(noise_array(tspace)) elem = tspace.element(arr_c, order=order) assert all_equal(elem, arr_c) assert elem.shape == elem.data.shape @@ -353,12 +373,14 @@ def test_element(tspace, odl_elem_order): if order is None or order == 'C': # None or same order should not lead to copy assert np.may_share_memory(elem.data, arr_c) + if order is not None: + assert _data_ptr(elem.data) == _data_ptr(arr_c) if order is not None: # Contiguousness in explicitly provided order should be guaranteed assert elem.data.flags[order + '_CONTIGUOUS'] - # From Numpy array (F order) - arr_f = np.asfortranarray(arr_c) + # From array (F order) + arr_f = _module(tspace.impl).asfortranarray(noise_array(tspace)) elem = tspace.element(arr_f, order=order) assert all_equal(elem, arr_f) assert elem.shape == elem.data.shape @@ -366,10 +388,19 @@ def test_element(tspace, odl_elem_order): if order is None or order == 'F': # None or same order should not lead to copy assert np.may_share_memory(elem.data, arr_f) + if order is not None: + assert _data_ptr(elem.data) == _data_ptr(arr_f) if order is not None: # Contiguousness in explicitly provided order should be guaranteed assert elem.data.flags[order + '_CONTIGUOUS'] + # From Numpy array + arr = np.random.rand(*tspace.shape).astype(tspace.dtype) + elem = tspace.element(arr, order=order) + assert all_equal(elem, arr) + assert elem.shape == elem.data.shape + assert elem.dtype == tspace.dtype == elem.data.dtype + # From pointer if tspace.impl == 'numpy': arr_c_ptr = arr_c.ctypes.data @@ -381,15 +412,14 @@ def test_element(tspace, odl_elem_order): assert all_equal(elem, arr_f) assert np.may_share_memory(elem.data, arr_f) - # Check errors - with pytest.raises(ValueError): - tspace.element(order='A') # only 'C' or 'F' valid + with pytest.raises(ValueError): + tspace.element(data_ptr=arr_c_ptr) # need order argument - with pytest.raises(ValueError): - tspace.element(data_ptr=arr_c_ptr) # need order argument + with pytest.raises(TypeError): + tspace.element(arr_c, arr_c_ptr) # forbidden to give both - with pytest.raises(TypeError): - tspace.element(arr_c, arr_c_ptr) # forbidden to give both + with pytest.raises(ValueError): + tspace.element(order='A') # only 'C', 'F' or None valid def test_equals_space(odl_tspace_impl): @@ -604,8 +634,8 @@ def test_multiply_exceptions(tspace): def test_power(tspace): """Test ``**`` against direct array exponentiation.""" [x_arr, y_arr], [x, y] = noise_elements(tspace, n=2) - y_pos = tspace.element(np.abs(y) + 0.1) - y_pos_arr = np.abs(y_arr) + 0.1 + y_pos = tspace.element(y.ufuncs.absolute() + 0.1) + y_pos_arr = (_module(tspace.impl).abs(y_arr) + 0.1).astype(tspace.dtype) # Testing standard positive integer power out-of-place and in-place assert all_almost_equal(x ** 2, x_arr ** 2) @@ -751,7 +781,7 @@ def test_norm(tspace): """Test the norm method against numpy.linalg.norm.""" xarr, x = noise_elements(tspace) - correct_norm = _module(tspace.impl).linalg.norm(xarr.ravel()) + correct_norm = np.linalg.norm(_as_numpy(xarr.ravel())) assert tspace.norm(x) == pytest.approx(correct_norm) assert x.norm() == pytest.approx(correct_norm) @@ -786,7 +816,7 @@ def test_dist(tspace): """Test the dist method against numpy.linalg.norm of the difference.""" [xarr, yarr], [x, y] = noise_elements(tspace, n=2) - correct_dist = _module(tspace.impl).linalg.norm((xarr - yarr).ravel()) + correct_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel())) assert tspace.dist(x, y) == pytest.approx(correct_dist) assert x.dist(y) == pytest.approx(correct_dist) @@ -817,6 +847,7 @@ def test_pdist(odl_tspace_impl, exponent): correct_dist = _module(impl).linalg.norm( (xarr - yarr).ravel(), ord=exponent) + assert space.dist(x, y) == pytest.approx(correct_dist) assert x.dist(y) == pytest.approx(correct_dist) @@ -891,7 +922,7 @@ def test_element_getitem_bool_array(odl_tspace_impl): x_arr_sliced = x_arr[cond_arr] x_sliced = x[cond] - assert all_equal(x_arr_sliced.asarray(), x_sliced.asarray()) + assert all_equal(x_arr_sliced, x_sliced) # Check that the space properties are preserved sliced_spc = x_sliced.space @@ -1216,13 +1247,7 @@ def test_array_weighting_inner(tspace): weighting_cls = _weighting_cls(tspace.impl, 'array') weighting = weighting_cls(weight_arr) - if tspace.impl == 'numpy': - true_inner = np.vdot(yarr, xarr * weight_arr) - elif tspace.impl == 'cupy': - true_inner = cupy.vdot(yarr, xarr * weight_arr) - else: - assert False - + true_inner = np.vdot(_as_numpy(yarr), _as_numpy(xarr * weight_arr)) assert weighting.inner(x, y) == pytest.approx(true_inner, rel=1e-2) # Exponent != 2 -> no inner product, should raise @@ -1238,18 +1263,12 @@ def test_array_weighting_norm(tspace, exponent): weighting_cls = _weighting_cls(tspace.impl, 'array') weighting = weighting_cls(weight_arr, exponent=exponent) - if tspace.impl == 'numpy': - norm = np.linalg.norm - elif tspace.impl == 'cupy': - norm = cupy.linalg.norm - else: - assert False - if exponent == float('inf'): - true_norm = float(norm((xarr).ravel(), ord=float('inf'))) + true_norm = np.linalg.norm(_as_numpy(xarr.ravel()), ord=float('inf')) else: - true_norm = float(norm((weight_arr ** (1 / exponent) * xarr).ravel(), - ord=exponent)) + true_norm = np.linalg.norm( + _as_numpy((weight_arr ** (1 / exponent) * xarr).ravel()), + ord=exponent) assert weighting.norm(x) == pytest.approx(true_norm, rel=1e-2) @@ -1263,12 +1282,11 @@ def test_array_weighting_dist(tspace, exponent): weighting = weighting_cls(weight_arr, exponent=exponent) if exponent == float('inf'): - true_dist = _module(tspace.impl).linalg.norm( - ((xarr - yarr)).ravel(), - ord=float('inf')) + true_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel()), + ord=float('inf')) else: - true_dist = _module(tspace.impl).linalg.norm( - (weight_arr ** (1 / exponent) * (xarr - yarr)).ravel(), + true_dist = np.linalg.norm( + _as_numpy((weight_arr ** (1 / exponent) * (xarr - yarr)).ravel()), ord=exponent) assert weighting.dist(x, y) == pytest.approx(true_dist, rel=1e-2) @@ -1335,16 +1353,16 @@ def test_const_weighting_inner(tspace): [xarr, yarr], [x, y] = noise_elements(tspace, 2) constant = 1.5 - true_result_const = constant * _module(tspace.impl).vdot(yarr, xarr) - weighting_cls = _weighting_cls(tspace.impl, 'const') - w_const = weighting_cls(constant) - assert w_const.inner(x, y) == pytest.approx(true_result_const, rel=1e-2) + weighting = weighting_cls(constant) + + true_inner = constant * np.vdot(_as_numpy(yarr), _as_numpy(xarr)) + assert weighting.inner(x, y) == pytest.approx(true_inner, rel=1e-2) # Exponent != 2 -> no inner - w_const = weighting_cls(constant, exponent=1) + weighing = weighting_cls(constant, exponent=1) with pytest.raises(NotImplementedError): - w_const.inner(x, y) + weighing.inner(x, y) def test_const_weighting_norm(tspace, exponent): @@ -1352,16 +1370,15 @@ def test_const_weighting_norm(tspace, exponent): xarr, x = noise_elements(tspace) constant = 1.5 + weighting_cls = _weighting_cls(tspace.impl, 'const') + weighting = weighting_cls(constant, exponent=exponent) + if exponent == float('inf'): factor = 1.0 else: factor = constant ** (1 / exponent) - true_norm = factor * _module(tspace.impl).linalg.norm( - xarr.ravel(), ord=exponent) - - weighting_cls = _weighting_cls(tspace.impl, 'const') - w_const = weighting_cls(constant, exponent=exponent) - assert w_const.norm(x) == pytest.approx(true_norm, rel=1e-2) + true_norm = factor * np.linalg.norm(_as_numpy(xarr.ravel()), ord=exponent) + assert weighting.norm(x) == pytest.approx(true_norm, rel=1e-2) def test_const_weighting_dist(tspace, exponent): @@ -1369,16 +1386,16 @@ def test_const_weighting_dist(tspace, exponent): [xarr, yarr], [x, y] = noise_elements(tspace, 2) constant = 1.5 + weighting_cls = _weighting_cls(tspace.impl, 'const') + weighting = weighting_cls(constant, exponent=exponent) + if exponent == float('inf'): factor = 1.0 else: factor = constant ** (1 / exponent) - true_dist = factor * _module(tspace.impl).linalg.norm( - (xarr - yarr).ravel(), ord=exponent) - - weighting_cls = _weighting_cls(tspace.impl, 'const') - w_const = weighting_cls(constant, exponent=exponent) - assert w_const.dist(x, y) == pytest.approx(true_dist, rel=1e-2) + true_dist = factor * np.linalg.norm(_as_numpy((xarr - yarr).ravel()), + ord=exponent) + assert weighting.dist(x, y) == pytest.approx(true_dist, rel=1e-2) def test_custom_inner(tspace): @@ -1386,7 +1403,7 @@ def test_custom_inner(tspace): [xarr, yarr], [x, y] = noise_elements(tspace, 2) def inner(x, y): - return _module(tspace.impl).vdot(y, x) + return np.vdot(_as_numpy(y.data), _as_numpy(x.data)) weighting_cls = _weighting_cls(tspace.impl, 'inner') w = weighting_cls(inner) @@ -1397,13 +1414,13 @@ def inner(x, y): assert w == w_same assert w != w_other - true_inner = inner(xarr, yarr) + true_inner = np.vdot(_as_numpy(yarr), _as_numpy(xarr)) assert w.inner(x, y) == pytest.approx(true_inner, rel=1e-2) - true_norm = _module(tspace.impl).linalg.norm(xarr.ravel()) + true_norm = np.linalg.norm(_as_numpy(xarr.ravel())) assert w.norm(x) == pytest.approx(true_norm, rel=1e-2) - true_dist = _module(tspace.impl).linalg.norm((xarr - yarr).ravel()) + true_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel())) assert w.dist(x, y) == pytest.approx(true_dist, rel=1e-2) with pytest.raises(TypeError): @@ -1414,10 +1431,11 @@ def test_custom_norm(tspace): """Test weighting with a custom norm.""" [xarr, yarr], [x, y] = noise_elements(tspace, 2) - norm = _module(tspace.impl).linalg.norm + def norm(x): + return np.linalg.norm(_as_numpy(x.data).ravel()) def other_norm(x): - return _module(tspace.impl).linalg.norm(x, ord=1) + return np.linalg.norm(_as_numpy(x.data).ravel(), ord=1) weighting_cls = _weighting_cls(tspace.impl, 'norm') w = weighting_cls(norm) @@ -1431,10 +1449,10 @@ def other_norm(x): with pytest.raises(NotImplementedError): w.inner(x, y) - true_norm = _module(tspace.impl).linalg.norm(xarr.ravel()) + true_norm = np.linalg.norm(_as_numpy(xarr.ravel())) assert w.norm(x) == pytest.approx(true_norm, rel=1e-2) - true_dist = _module(tspace.impl).linalg.norm((xarr - yarr).ravel()) + true_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel())) assert w.dist(x, y) == pytest.approx(true_dist, rel=1e-2) with pytest.raises(TypeError): @@ -1446,10 +1464,10 @@ def test_custom_dist(tspace): [xarr, yarr], [x, y] = noise_elements(tspace, 2) def dist(x, y): - return _module(tspace.impl).linalg.norm(x - y) + return np.linalg.norm(_as_numpy((x - y).data).ravel()) def other_dist(x, y): - return _module(tspace.impl).linalg.norm(x - y, ord=1) + return np.linalg.norm(_as_numpy((x - y).data).ravel(), ord=1) weighting_cls = _weighting_cls(tspace.impl, 'dist') w = weighting_cls(dist) @@ -1466,7 +1484,7 @@ def other_dist(x, y): with pytest.raises(NotImplementedError): w.norm(x) - true_dist = _module(tspace.impl).linalg.norm((xarr - yarr).ravel()) + true_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel())) assert w.dist(x, y) == pytest.approx(true_dist, rel=1e-2) with pytest.raises(TypeError): From d456e6ec6f63e2c2f594ee0485a2cb9786a1c9c9 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Mon, 27 Nov 2017 13:56:06 +0100 Subject: [PATCH 18/38] WIP: fix cupy tests --- odl/space/cupy_tensors.py | 40 +++++- odl/test/space/tensors_test.py | 225 ++++++++++++++++++++------------- 2 files changed, 171 insertions(+), 94 deletions(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index 1195ac726c7..c4616046b81 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -1613,8 +1613,12 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): input1, ..., inputN : Positional arguments to ``ufunc.method``. + force_native : bool, optional + If ``True``, raise a ``ValueError`` if there is no native CuPy + function available for the given ``ufunc``, ``method`` and inputs. + Default : ``False`` kwargs : - Keyword arguments to ``ufunc.method``. + Further keyword arguments passed to ``ufunc.method``. Returns ------- @@ -1788,6 +1792,13 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): use_native = (native_ufunc is not None and hasattr(native_ufunc, method)) + force_native = kwargs.pop('force_native', False) + if force_native and not use_native: + raise ValueError( + 'no native function available to evaluate `{}.{}` with ' + 'inputs {!r}, kwargs {!r} and `out` {!r}' + ''.format(ufunc.__name__, method, inputs, kwargs, out_tuple)) + # Assign to `out` or `out1` and `out2`, respectively, unwrapping the # data container out = out1 = out2 = None @@ -1916,18 +1927,30 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): use_native = (use_native and native_method is not None) def eval_at_via_npy(*inputs, **kwargs): + """Use Numpy to evaluate ufunc.at, converting the inputs.""" import ctypes - cupy_arr = inputs[0] - npy_arr = cupy.asnumpy(cupy_arr) - new_inputs = (npy_arr,) + inputs[1:] + + if force_native: + raise ValueError( + 'no native function available to evaluate `{}.at` ' + 'with inputs {!r} and kwargs {!r}' + ''.format(ufunc.__name__, inputs, kwargs)) + new_inputs = [cupy.asnumpy(inputs[0]), inputs[1]] + if ufunc.nin == 2: + new_inputs.append(cupy.asnumpy(inputs[2])) + + # Shouldn't exist, but let upstream code raise + new_inputs.extend(inputs[3:]) + super(CupyTensor, self).__array_ufunc__( ufunc, method, *new_inputs, **kwargs) # Workaround for assignment cupy_arr[:] = npy_arr not # working. See # https://github.com/cupy/cupy/issues/593 and # https://github.com/odlgroup/odl/issues/1248 - cupy_arr.data.copy_from_host( - npy_arr.ctypes.data_as(ctypes.c_void_p), npy_arr.nbytes) + inputs[0].data.copy_from_host( + new_inputs[0].ctypes.data_as(ctypes.c_void_p), + new_inputs[0].nbytes) if use_native: # Native method could exist but raise `NotImplementedError` @@ -1952,6 +1975,11 @@ def eval_at_via_npy(*inputs, **kwargs): else: function = cupy.prod if method == 'reduce' else cupy.cumprod + # Make 0 the default for `axis` as the ufunc methods. The default + # for `[cum]sum` and `[cum]prod` is to reduce over all axes + axis = kwargs.pop('axis', 0) + kwargs['axis'] = axis + kwargs['out'] = out res = function(*inputs, **kwargs) # Shortcut for scalar return value diff --git a/odl/test/space/tensors_test.py b/odl/test/space/tensors_test.py index b884e5cc0f0..819c18b92ff 100644 --- a/odl/test/space/tensors_test.py +++ b/odl/test/space/tensors_test.py @@ -1499,9 +1499,27 @@ def testodl_ufuncs(tspace, odl_ufunc): name = odl_ufunc # Get the ufunc from numpy as reference, plus some additional info - npy_ufunc = getattr(np, name) - nin = npy_ufunc.nin - nout = npy_ufunc.nout + ufunc_npy = getattr(np, name) + nin = ufunc_npy.nin + nout = ufunc_npy.nout + + def _check_result_type(result, expected_type): + if nout == 1: + assert isinstance(result, tspace.element_type) + elif nout > 1: + for i in range(nout): + assert isinstance(result[i], tspace.element_type) + else: + assert False + + def _check_result_is_out(result, out_seq): + if nout == 1: + assert result is out_seq[0] + elif nout > 1: + for i in range(nout): + assert result[i] is out_seq[i] + else: + assert False if (np.issubsctype(tspace.dtype, np.floating) or np.issubsctype(tspace.dtype, np.complexfloating) and @@ -1538,10 +1556,10 @@ def testodl_ufuncs(tspace, odl_ufunc): # Create some data arrays, elements = noise_elements(tspace, nin + nout) - in_arrays = arrays[:nin] - out_arrays = arrays[nin:] + # Arrays of the space's own data storage type + in_arrays_own = arrays[:nin] + in_arrays_npy = [_as_numpy(arr) for arr in arrays[:nin]] data_elem = elements[0] - out_elems = elements[nin:] if nout == 1: out_arr_kwargs = {'out': out_arrays[0]} @@ -1550,116 +1568,147 @@ def testodl_ufuncs(tspace, odl_ufunc): out_arr_kwargs = {'out': out_arrays[:nout]} out_elem_kwargs = {'out': out_elems[:nout]} + # Get function to call, using both interfaces: - # - vec.ufunc(other_args) - # - np.ufunc(vec, other_args) - elem_fun_old = getattr(data_elem.ufuncs, name) - in_elems_old = elements[1:nin] - elem_fun_new = npy_ufunc - in_elems_new = elements[:nin] - - # Out-of-place - npy_result = npy_ufunc(*in_arrays) - odl_result_old = elem_fun_old(*in_elems_old) - assert all_almost_equal(npy_result, odl_result_old) - odl_result_new = elem_fun_new(*in_elems_new) - assert all_almost_equal(npy_result, odl_result_new) - - # Test type of output - if nout == 1: - assert isinstance(odl_result_old, tspace.element_type) - assert isinstance(odl_result_new, tspace.element_type) - elif nout > 1: - for i in range(nout): - assert isinstance(odl_result_old[i], tspace.element_type) - assert isinstance(odl_result_new[i], tspace.element_type) - - # In-place with ODL objects as `out` - npy_result = npy_ufunc(*in_arrays, **out_arr_kwargs) - odl_result_old = elem_fun_old(*in_elems_old, **out_elem_kwargs) - assert all_almost_equal(npy_result, odl_result_old) + # - vec.ufunc(*other_args) + # - np.ufunc(vec, *other_args) + ufunc_method = getattr(data_elem.ufuncs, name) + in_elems_method = elements[1:nin] + in_elems_npy = elements[:nin] + result_npy = ufunc_npy(*in_arrays_npy) + + # Out-of-place -- in = elements -- ufunc = method or numpy + result = ufunc_method(*in_elems_method) + assert all_almost_equal(result_npy, result) + _check_result_type(result, tspace.element_type) + + result = ufunc_npy(*in_elems_npy) + assert all_almost_equal(result_npy, result) + _check_result_type(result, tspace.element_type) + + # Out-of-place -- in = numpy or own arrays -- ufunc = method + result = ufunc_method(*in_arrays_npy[1:]) + assert all_almost_equal(result_npy, result) + _check_result_type(result, tspace.element_type) + + result = ufunc_method(*in_arrays_own[1:]) + assert all_almost_equal(result_npy, result) + _check_result_type(result, tspace.element_type) + + # In-place -- in = elements -- out = elements -- ufunc = method or numpy + result = ufunc_method(*in_elems_method, **out_elem_kwargs) + assert all_almost_equal(result_npy, result) + _check_result_is_out(result, out_elems[:nout]) if USE_ARRAY_UFUNCS_INTERFACE: - # In-place will not work with Numpy < 1.13 - odl_result_new = elem_fun_new(*in_elems_new, **out_elem_kwargs) - assert all_almost_equal(npy_result, odl_result_new) + # Custom objects not allowed as `out` for numpy < 1.13 + result = ufunc_npy(*in_elems_npy, **out_elem_kwargs) + assert all_almost_equal(result_npy, result) + _check_result_is_out(result, out_elems[:nout]) - # Check that returned stuff refers to given out - if nout == 1: - assert odl_result_old is out_elems[0] - if USE_ARRAY_UFUNCS_INTERFACE: - assert odl_result_new is out_elems[0] - elif nout > 1: - for i in range(nout): - assert odl_result_old[i] is out_elems[i] - if USE_ARRAY_UFUNCS_INTERFACE: - assert odl_result_new[i] is out_elems[i] - - # In-place with Numpy array as `out` for new interface + # In-place -- in = elements -- out = numpy or own arrays -- ufunc = numpy + # This case is only supported with the new interface if USE_ARRAY_UFUNCS_INTERFACE: - out_arrays_new = [np.empty_like(arr) for arr in out_arrays] + # Fresh arrays for output + out_arrays_npy = [np.empty_like(_as_numpy(arr)) + for arr in arrays[nin:]] + out_arrays_own = [_module(tspace.impl).empty_like(arr) + for arr in arrays[nin:]] if nout == 1: - out_elem_kwargs_new = {'out': out_arrays_new[0]} + kwargs_npy = {'out': out_arrays_npy[0]} + kwargs_own = {'out': out_arrays_own[0]} elif nout > 1: - out_elem_kwargs_new = {'out': out_arrays_new[:nout]} - - odl_result_elem_new = elem_fun_new(*in_elems_new, - **out_elem_kwargs_new) - assert all_almost_equal(npy_result, odl_result_elem_new) + kwargs_npy = {'out': out_arrays_npy[:nout]} + kwargs_own = {'out': out_arrays_own[:nout]} - if nout == 1: - assert odl_result_elem_new is out_arrays_new[0] - elif nout > 1: - for i in range(nout): - assert odl_result_elem_new[i] is out_arrays_new[i] + result_out_npy = ufunc_npy(*in_elems_npy, **kwargs_npy) + result_out_own = ufunc_npy(*in_elems_npy, **kwargs_own) + assert all_almost_equal(result_out_npy, result_npy) + assert all_almost_equal(result_out_own, result_npy) + _check_result_is_out(result_out_npy, out_arrays_npy) + _check_result_is_out(result_out_own, out_arrays_own) if USE_ARRAY_UFUNCS_INTERFACE: # Check `ufunc.at` indices = [[0, 0, 1], [0, 1, 2]] - mod_array = in_arrays[0].copy() - mod_elem = in_elems_new[0].copy() + mod_array = in_arrays_npy[0].copy() + mod_elem = in_elems_npy[0].copy() if nin == 1: - npy_result = npy_ufunc.at(mod_array, indices) - odl_result = npy_ufunc.at(mod_elem, indices) + result_npy = ufunc_npy.at(mod_array, indices) + result = ufunc_npy.at(mod_elem, indices) elif nin == 2: - other_array = in_arrays[1][indices] - other_elem = in_elems_new[1][indices] - npy_result = npy_ufunc.at(mod_array, indices, other_array) - odl_result = npy_ufunc.at(mod_elem, indices, other_elem) + other_array = in_arrays_npy[1][indices] + other_elem = in_elems_npy[1][indices] + result_npy = ufunc_npy.at(mod_array, indices, other_array) + result = ufunc_npy.at(mod_elem, indices, other_elem) - assert all_almost_equal(odl_result, npy_result) + assert all_almost_equal(result, result_npy) # Check `ufunc.reduce` if nin == 2 and nout == 1 and USE_ARRAY_UFUNCS_INTERFACE: - in_array = in_arrays[0] - in_elem = in_elems_new[0] + in_array = in_arrays_npy[0] + in_elem = in_elems_npy[0] # We only test along one axis since some binary ufuncs are not # re-orderable, in which case Numpy raises a ValueError - npy_result = npy_ufunc.reduce(in_array) - odl_result = npy_ufunc.reduce(in_elem) - assert all_almost_equal(odl_result, npy_result) - odl_result_keepdims = npy_ufunc.reduce(in_elem, keepdims=True) - assert odl_result_keepdims.shape == (1,) + in_elem.shape[1:] - # In-place using `out` (with ODL vector and array) - out_elem = odl_result_keepdims.space.element() - out_array = np.empty(odl_result_keepdims.shape, - dtype=odl_result_keepdims.dtype) - npy_ufunc.reduce(in_elem, out=out_elem, keepdims=True) - npy_ufunc.reduce(in_elem, out=out_array, keepdims=True) - assert all_almost_equal(out_elem, odl_result_keepdims) - assert all_almost_equal(out_array, odl_result_keepdims) + + # Out-of-place -- in = element + result_npy = ufunc_npy.reduce(in_array) + result = ufunc_npy.reduce(in_elem) + assert all_almost_equal(result, result_npy) + result_keepdims = ufunc_npy.reduce(in_elem, keepdims=True) + assert result_keepdims.shape == (1,) + in_elem.shape[1:] + + # In-place -- in = element -- out = element or numpy array or own array + out_elem = result_keepdims.space.element() + ufunc_npy.reduce(in_elem, out=out_elem, keepdims=True) + assert all_almost_equal(out_elem, result_keepdims) + out_array_npy = np.empty(result_keepdims.shape, + dtype=result_keepdims.dtype) + ufunc_npy.reduce(in_elem, out=out_array_npy, keepdims=True) + assert all_almost_equal(out_array_npy, result_keepdims) + out_array_own = _module(tspace.impl).empty( + result_keepdims.shape, dtype=result_keepdims.dtype) + ufunc_npy.reduce(in_elem, out=out_array_own, keepdims=True) + assert all_almost_equal(out_array_own, result_keepdims) + # Using a specific dtype - npy_result = npy_ufunc.reduce(in_array, dtype=complex) - odl_result = npy_ufunc.reduce(in_elem, dtype=complex) - assert odl_result.dtype == npy_result.dtype - assert all_almost_equal(odl_result, npy_result) + result_npy = ufunc_npy.reduce(in_array, dtype=complex) + result = ufunc_npy.reduce(in_elem, dtype=complex) + assert result.dtype == result_npy.dtype + assert all_almost_equal(result, result_npy) # Other ufunc method use the same interface, to we don't perform # extra tests for them. +def test_ufunc_cupy_force_native(): + """Test the ``force_native`` flag for cupy based ufuncs.""" + if not USE_ARRAY_UFUNCS_INTERFACE: + pytest.skip('`force_native` option only used in __array_ufuncs__') + + space = odl.rn((3, 4), impl='cupy') + + # Make sure we call native code for supported ufuncs + for ufunc in [np.sin, np.absolute, np.add, np.remainder, np.fmod]: + nin, nout = ufunc.nin, ufunc.nout + _, in_elems = noise_elements(space, n=2) + out_arrays, out_elems = noise_elements(space, n=2) + ufunc(*in_elems[:nin], out=out_elems[:nout], force_native=True) + ufunc(*in_elems[:nin], out=out_arrays[:nout], force_native=True) + + # These have explicit native implementations + for ufunc in [np.add, np.multiply]: + for method in ['reduce', 'accumulate']: + in_elem = noise_element(space) + getattr(ufunc, method)(in_elem, force_native=True) + + for ufunc in [np.minimum, np.maximum]: + in_elem = noise_element(space) + ufunc.reduce(in_elem, force_native=True) + + def test_ufunc_corner_cases(odl_tspace_impl): """Check if some corner cases are handled correctly.""" impl = odl_tspace_impl From f0136b57772d23b234745cc2c3aa0deac98d2b04 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Mon, 27 Nov 2017 18:52:13 +0100 Subject: [PATCH 19/38] ENH: make ufuncs and ufunc_ops work with 2 outputs --- odl/ufunc_ops/ufunc_ops.py | 43 ++++++------ odl/util/ufuncs.py | 133 ++++++++++++++++++++++++------------- 2 files changed, 109 insertions(+), 67 deletions(-) diff --git a/odl/ufunc_ops/ufunc_ops.py b/odl/ufunc_ops/ufunc_ops.py index af88b1f408a..2e66b711240 100644 --- a/odl/ufunc_ops/ufunc_ops.py +++ b/odl/ufunc_ops/ufunc_ops.py @@ -209,10 +209,10 @@ def derivative(self, point): return derivative -def ufunc_class_factory(name, nargin, nargout, docstring): +def ufunc_class_factory(name, nin, nout, docstring): """Create a Ufunc `Operator` from a given specification.""" - assert 0 <= nargin <= 2 + assert 0 <= nin <= 2 def __init__(self, space): """Initialize an instance. @@ -225,21 +225,22 @@ def __init__(self, space): if not isinstance(space, LinearSpace): raise TypeError('`space` {!r} not a `LinearSpace`'.format(space)) - if nargin == 1: + if nin == 1: domain = space0 = space dtypes = [space.dtype] - elif nargin == len(space) == 2 and isinstance(space, ProductSpace): + elif nin == len(space) == 2 and isinstance(space, ProductSpace): domain = space space0 = space[0] dtypes = [space[0].dtype, space[1].dtype] else: - domain = ProductSpace(space, nargin) + domain = ProductSpace(space, nin) space0 = space dtypes = [space.dtype, space.dtype] dts_out = dtypes_out(name, dtypes) + print(dts_out) - if nargout == 1: + if nout == 1: range = space0.astype(dts_out[0]) else: range = ProductSpace(space0.astype(dts_out[0]), @@ -253,12 +254,12 @@ def _call(self, x, out=None): # TODO: use `__array_ufunc__` when implemented on `ProductSpace`, # or try both if out is None: - if nargin == 1: + if nin == 1: return getattr(x.ufuncs, name)() else: return getattr(x[0].ufuncs, name)(*x[1:]) else: - if nargin == 1: + if nin == 1: return getattr(x.ufuncs, name)(out=out) else: return getattr(x[0].ufuncs, name)(*x[1:], out=out) @@ -274,19 +275,19 @@ def __repr__(self): dtype = float space = tensor_space(3, dtype=dtype) - if nargin == 1: + if nin == 1: vec = space.element([-1, 1, 2]) arg = '{}'.format(vec) with np.errstate(all='ignore'): result = getattr(vec.ufuncs, name)() else: vec = space.element([-1, 1, 2]) - vec2 = space.element([3, 4, 5]) + vec2 = [3, 4, 5] arg = '[{}, {}]'.format(vec, vec2) with np.errstate(all='ignore'): result = getattr(vec.ufuncs, name)(vec2) - if nargout == 2: + if nout == 2: result_space = ProductSpace(vec.space, 2) result = repr(result_space.element(result)) @@ -305,10 +306,10 @@ def __repr__(self): return type(full_name, (Operator,), attributes) -def ufunc_functional_factory(name, nargin, nargout, docstring): +def ufunc_functional_factory(name, nin, nout, docstring): """Create a ufunc `Functional` from a given specification.""" - assert 0 <= nargin <= 2 + assert 0 <= nin <= 2 def __init__(self, field): """Initialize an instance. @@ -330,7 +331,7 @@ def __init__(self, field): def _call(self, x): """Return ``self(x)``.""" - if nargin == 1: + if nin == 1: return getattr(np, name)(x) else: return getattr(np, name)(*x) @@ -341,10 +342,10 @@ def __repr__(self): # Create example (also functions as doctest) - if nargin != 1: + if nin != 1: raise NotImplementedError('Currently not suppored') - if nargout != 1: + if nout != 1: raise NotImplementedError('Currently not suppored') space = RealNumbers() @@ -394,7 +395,7 @@ def __repr__(self): # Create an operator for each ufunc -for name, nargin, nargout, docstring in UFUNCS: +for name, nin, nout, docstring in UFUNCS: def indirection(name, docstring): # Indirection is needed since name should be saved but is changed # in the loop. @@ -410,17 +411,17 @@ def ufunc_factory(domain=RealNumbers()): raise ValueError('ufunc not available for {}'.format(domain)) return ufunc_factory - globals()[name + '_op'] = ufunc_class_factory(name, nargin, - nargout, docstring) + globals()[name + '_op'] = ufunc_class_factory(name, nin, + nout, docstring) if not _is_integer_only_ufunc(name): operator_example = RAW_UFUNC_FACTORY_OPERATOR_DOCSTRING.format( name=name) else: operator_example = "" - if not _is_integer_only_ufunc(name) and nargin == 1 and nargout == 1: + if not _is_integer_only_ufunc(name) and nin == 1 and nout == 1: globals()[name + '_func'] = ufunc_functional_factory( - name, nargin, nargout, docstring) + name, nin, nout, docstring) functional_example = RAW_UFUNC_FACTORY_FUNCTIONAL_DOCSTRING.format( name=name) else: diff --git a/odl/util/ufuncs.py b/odl/util/ufuncs.py index caa6f1fbb0d..d2694553140 100644 --- a/odl/util/ufuncs.py +++ b/odl/util/ufuncs.py @@ -32,25 +32,35 @@ __all__ = ('TensorSpaceUfuncs', 'ProductSpaceUfuncs') -# Some are ignored since they don't cooperate with dtypes, needs fix -RAW_UFUNCS = ['absolute', 'add', 'arccos', 'arccosh', 'arcsin', 'arcsinh', - 'arctan', 'arctan2', 'arctanh', 'bitwise_and', 'bitwise_or', - 'bitwise_xor', 'ceil', 'conj', 'copysign', 'cos', 'cosh', - 'deg2rad', 'divide', 'equal', 'exp', 'exp2', 'expm1', 'floor', - 'floor_divide', 'fmax', 'fmin', 'fmod', 'greater', - 'greater_equal', 'hypot', 'invert', 'isfinite', 'isinf', 'isnan', - 'left_shift', 'less', 'less_equal', 'log', 'log10', 'log1p', - 'log2', 'logaddexp', 'logaddexp2', 'logical_and', 'logical_not', - 'logical_or', 'logical_xor', 'maximum', 'minimum', 'mod', 'modf', - 'multiply', 'negative', 'not_equal', 'power', - 'rad2deg', 'reciprocal', 'remainder', 'right_shift', 'rint', - 'sign', 'signbit', 'sin', 'sinh', 'sqrt', 'square', 'subtract', - 'tan', 'tanh', 'true_divide', 'trunc'] -# ,'isreal', 'iscomplex', 'ldexp', 'frexp' +_npy_maj, _npy_min = [int(n) for n in np.__version__.split('.')[:2]] + +# Supported by Numpy 1.9 and higher +UFUNC_NAMES = [ + 'absolute', 'add', 'arccos', 'arccosh', 'arcsin', 'arcsinh', 'arctan', + 'arctan2', 'arctanh', 'bitwise_and', 'bitwise_or', 'bitwise_xor', 'ceil', + 'conj', 'conjugate', 'copysign', 'cos', 'cosh', 'deg2rad', 'degrees', + 'divide', 'equal', 'exp', 'exp2', 'expm1', 'fabs', 'floor', 'floor_divide', + 'fmax', 'fmin', 'fmod', 'frexp', 'greater', 'greater_equal', 'hypot', + 'invert', 'isfinite', 'isinf', 'isnan', 'ldexp', 'left_shift', 'less', + 'less_equal', 'log', 'log10', 'log1p', 'log2', 'logaddexp', 'logaddexp2', + 'logical_and', 'logical_not', 'logical_or', 'logical_xor', 'maximum', + 'minimum', 'mod', 'modf', 'multiply', 'negative', 'nextafter', + 'not_equal', 'power', 'rad2deg', 'radians', 'reciprocal', 'remainder', + 'right_shift', 'rint', 'sign', 'signbit', 'sin', 'sinh', 'sqrt', + 'square', 'spacing', 'subtract', 'tan', 'tanh', 'true_divide', 'trunc'] + +if (_npy_maj, _npy_min) >= (1, 10): + UFUNC_NAMES.extend(['abs', 'cbrt', 'bitwise_not']) + +if (_npy_maj, _npy_min) >= (1, 12): + UFUNC_NAMES.extend(['float_power']) + +if (_npy_maj, _npy_min) >= (1, 13): + UFUNC_NAMES.extend(['divmod', 'heaviside', 'positive']) # Add some standardized information UFUNCS = [] -for name in RAW_UFUNCS: +for name in UFUNC_NAMES: ufunc = getattr(np, name) n_in, n_out = ufunc.nin, ufunc.nout descr = ufunc.__doc__.splitlines()[2] @@ -102,6 +112,14 @@ def wrapper(self, x2, out=None, **kwargs): return self.elem.__array_ufunc__( ufunc, '__call__', self.elem, x2, out=(out,), **kwargs) + elif n_out == 2: + def wrapper(self, x2, out=None, **kwargs): + if out is None: + out = (None, None) + + return self.elem.__array_ufunc__( + ufunc, '__call__', self.elem, x2, out=out, **kwargs) + else: raise NotImplementedError else: @@ -188,24 +206,35 @@ def wrap_ufunc_productspace(name, n_in, n_out, doc): if n_in == 1: if n_out == 1: def wrapper(self, out=None, **kwargs): + from odl.space.pspace import ProductSpace if out is None: - result = [getattr(x.ufuncs, name)(**kwargs) - for x in self.elem] - return self.elem.space.element(result) - else: - for x, out_x in zip(self.elem, out): - getattr(x.ufuncs, name)(out=out_x, **kwargs) - return out + out = [None] * len(self.elem.space) + + res = [] + for xi, out_i in zip(self.elem, out): + r = getattr(xi.ufuncs, name)(out=out_i, **kwargs) + res.append(r) + out_space = ProductSpace(*[r.space for r in res]) + return out_space.element(res) elif n_out == 2: def wrapper(self, out1=None, out2=None, **kwargs): + from odl.space.pspace import ProductSpace if out1 is None: - out1 = self.elem.space.element() + out1 = [None] * len(self.elem.space) if out2 is None: - out2 = self.elem.space.element() - for x, out1_x, out2_x in zip(self.elem, out1, out2): - getattr(x.ufuncs, name)(out1=out1_x, out2=out2_x, **kwargs) - return out1, out2 + out2 = [None] * len(self.elem.space) + + res1, res2 = [], [] + for xi, out1_i, out2_i in zip(self.elem, out1, out2): + r1, r2 = getattr(xi.ufuncs, name)(out1=out1_i, + out2=out2_i, + **kwargs) + res1.append(r1) + res2.append(r2) + out_space_1 = ProductSpace(*[r.space for r in res1]) + out_space_2 = ProductSpace(*[r.space for r in res2]) + return out_space_1.element(res1), out_space_2.element(res2) else: raise NotImplementedError @@ -213,24 +242,36 @@ def wrapper(self, out1=None, out2=None, **kwargs): elif n_in == 2: if n_out == 1: def wrapper(self, x2, out=None, **kwargs): - if x2 in self.elem.space: - if out is None: - result = [getattr(x.ufuncs, name)(x2p, **kwargs) - for x, x2p in zip(self.elem, x2)] - return self.elem.space.element(result) - else: - for x, x2p, outp in zip(self.elem, x2, out): - getattr(x.ufuncs, name)(x2p, out=outp, **kwargs) - return out - else: - if out is None: - result = [getattr(x.ufuncs, name)(x2, **kwargs) - for x in self.elem] - return self.elem.space.element(result) - else: - for x, outp in zip(self.elem, out): - getattr(x.ufuncs, name)(x2, out=outp, **kwargs) - return out + from odl.space.pspace import ProductSpace + if out is None: + out = [None] * len(self.elem.space) + + res = [] + for x1_i, x2_i, out_i in zip(self.elem, x2, out): + r = getattr(x1_i.ufuncs, name)(x2_i, out=out_i, **kwargs) + res.append(r) + out_space = ProductSpace(*[r.space for r in res]) + return out_space.element(res) + + elif n_out == 2: + def wrapper(self, x2, out1=None, out2=None, **kwargs): + from odl.space.pspace import ProductSpace + if out1 is None: + out1 = [None] * len(self.elem.space) + if out2 is None: + out2 = [None] * len(self.elem.space) + + res1, res2 = [], [] + for x1_i, x2_i, out1_i, out2_i in zip(self.elem, x2, out1, + out2): + r1, r2 = getattr(x1_i.ufuncs, name)(x2_i, out1=out1_i, + out2=out2_i, + **kwargs) + res1.append(r1) + res2.append(r2) + out_space_1 = ProductSpace(*[r.space for r in res1]) + out_space_2 = ProductSpace(*[r.space for r in res2]) + return out_space_1.element(res1), out_space_2.element(res2) else: raise NotImplementedError From 1e8942593a873612bf28c93f99ab12f60614d49d Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Mon, 27 Nov 2017 18:52:41 +0100 Subject: [PATCH 20/38] MAINT: make weighting methods numpy-independent --- odl/space/weighting.py | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/odl/space/weighting.py b/odl/space/weighting.py index 8fcabbf50d8..8e051c9e169 100644 --- a/odl/space/weighting.py +++ b/odl/space/weighting.py @@ -500,7 +500,7 @@ def array(self): def is_valid(self): """Return True if the array is a valid weight, i.e. positive.""" - return np.all(np.greater(self.array, 0)) + return (self.array > 0).all() def __eq__(self, other): """Return ``self == other``. @@ -547,9 +547,9 @@ def equiv(self, other): elif isinstance(other, MatrixWeighting): return other.equiv(self) elif isinstance(other, ConstWeighting): - return np.array_equiv(self.array, other.const) + return (self.array == other.const).all() else: - return np.array_equal(self.array, other.array) + return (self.array == other.array).all() @property def repr_part(self): From 41fe853614d2fcdab02701d5fa69d07e041b9ae3 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Mon, 27 Nov 2017 18:52:53 +0100 Subject: [PATCH 21/38] WIP: fix cupy tests --- odl/space/cupy_tensors.py | 74 ++++++++++++++++++++++------------ odl/test/space/tensors_test.py | 18 +++++---- 2 files changed, 58 insertions(+), 34 deletions(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index c4616046b81..60f9df5131b 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -608,13 +608,14 @@ def __init__(self, shape, dtype='float64', device=None, **kwargs): norm = kwargs.pop('norm', None) inner = kwargs.pop('inner', None) weighting = kwargs.pop('weighting', None) - exponent = kwargs.pop('exponent', 2.0) + exponent = kwargs.pop('exponent', None) # Check validity of option combination (3 or 4 out of 4 must be None) if sum(x is None for x in (dist, norm, inner, weighting)) < 3: raise ValueError('invalid combination of options `weighting`, ' '`dist`, `norm` and `inner`') - if any(x is not None for x in (dist, norm, inner)) and exponent != 2.0: + if (any(x is not None for x in (dist, norm, inner)) and + exponent is not None): raise ValueError('`exponent` cannot be used together with ' '`dist`, `norm` and `inner`') @@ -624,7 +625,7 @@ def __init__(self, shape, dtype='float64', device=None, **kwargs): if weighting.impl != 'cupy': raise ValueError("`weighting.impl` must be 'cupy', " '`got {!r}'.format(weighting.impl)) - if weighting.exponent != exponent: + if exponent is not None and weighting.exponent != exponent: raise ValueError('`weighting.exponent` conflicts with ' '`exponent`: {} != {}' ''.format(weighting.exponent, exponent)) @@ -653,7 +654,7 @@ def __init__(self, shape, dtype='float64', device=None, **kwargs): elif inner is not None: self.__weighting = CupyTensorSpaceCustomInner(inner) else: # all None -> no weighing - self.__weighting = CupyTensorSpaceConstWeighting(1.0, exponent) + self.__weighting = _weighting(1.0, exponent) @property def device(self): @@ -1400,15 +1401,18 @@ def __getitem__(self, indices): else: weighting = None + kwargs = {} if isinstance(weighting, CupyTensorSpaceArrayWeighting): weighting = weighting.array[indices] elif isinstance(weighting, CupyTensorSpaceConstWeighting): # Axes were removed, cannot infer new constant if arr.ndim != self.ndim: weighting = None + kwargs['exponent'] = self.space.exponent space = type(self.space)(arr.shape, dtype=self.dtype, - device=self.device, weighting=weighting) + device=self.device, weighting=weighting, + **kwargs) return space.element(arr) def __setitem__(self, indices, values): @@ -1770,6 +1774,12 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): "need 0 or 1 `out` arguments for `method={!r}`, " 'got {}'.format(method, len(out_tuple))) + # Catch wrong number of inputs + if method == '__call__' and len(inputs) != ufunc.nin: + raise ValueError( + "need {} inputs for `method='__call__'`, got {}" + ''.format(ufunc.nin, len(inputs))) + # We allow our own tensors, the data container type and # `numpy.ndarray` objects as `out` (see docs for reason for the # latter) @@ -1778,27 +1788,6 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): for o in out_tuple): return NotImplemented - # Determine native ufunc vs. Numpy ufunc - if any(isinstance(o, np.ndarray) for o in out_tuple): - native_ufunc = None - use_native = False - else: - native_ufunc = getattr(cupy, ufunc.__name__, None) - # Manual assignment for sum, cumsum, prod and cumprod - if (ufunc in (np.add, np.multiply) and - method in ('reduce', 'accumulate')): - use_native = native_ufunc is not None - else: - use_native = (native_ufunc is not None and - hasattr(native_ufunc, method)) - - force_native = kwargs.pop('force_native', False) - if force_native and not use_native: - raise ValueError( - 'no native function available to evaluate `{}.{}` with ' - 'inputs {!r}, kwargs {!r} and `out` {!r}' - ''.format(ufunc.__name__, method, inputs, kwargs, out_tuple)) - # Assign to `out` or `out1` and `out2`, respectively, unwrapping the # data container out = out1 = out2 = None @@ -1824,6 +1813,30 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): inp.data if isinstance(inp, type(self)) else inp for inp in inputs) + # Determine native ufunc vs. Numpy ufunc + if any(isinstance(o, np.ndarray) for o in out_tuple): + native_ufunc = None + use_native = False + elif any(not isinstance(inp, type(self.data)) for inp in inputs): + native_ufunc = None + use_native = False + else: + native_ufunc = getattr(cupy, ufunc.__name__, None) + # Manual assignment for sum, cumsum, prod and cumprod + if (ufunc in (np.add, np.multiply) and + method in ('reduce', 'accumulate')): + use_native = native_ufunc is not None + else: + use_native = (native_ufunc is not None and + hasattr(native_ufunc, method)) + + force_native = kwargs.pop('force_native', False) + if force_native and not use_native: + raise ValueError( + 'no native function available to evaluate `{}.{}` with ' + 'inputs {!r}, kwargs {!r} and `out` {!r}' + ''.format(ufunc.__name__, method, inputs, kwargs, out_tuple)) + if use_native: # TODO: remove when upstream issue is fixed # For native ufuncs, we turn non-scalar inputs into cupy arrays @@ -2200,6 +2213,9 @@ def __ipow__(self, other): def _weighting(weights, exponent): """Return a weighting whose type is inferred from the arguments.""" + if exponent is None: + exponent = 2.0 + if np.isscalar(weights): weighting = CupyTensorSpaceConstWeighting(weights, exponent=exponent) else: @@ -2498,6 +2514,12 @@ def dist(self, x1, x2): return float(distpw(x1.data, x2.data, self.exponent, self.array, out)) + def __hash__(self): + """Return ``hash(self)``.""" + # CuPy array has no `tobytes` + return hash((super(ArrayWeighting, self).__hash__(), + cupy.asnumpy(self.array).tobytes())) + # TODO: remove repr_part and __repr__ when cupy.ndarray.__array__ # is implemented. See # https://github.com/cupy/cupy/issues/589 diff --git a/odl/test/space/tensors_test.py b/odl/test/space/tensors_test.py index 819c18b92ff..a6567cf3d34 100644 --- a/odl/test/space/tensors_test.py +++ b/odl/test/space/tensors_test.py @@ -898,13 +898,14 @@ def test_element_setitem(odl_tspace_impl, setitem_indices): assert all_equal(x, x_arr) # Setting values with arrays - rhs_arr = np.ones(sliced_shape) + rhs_arr = _module(tspace_impl).ones(sliced_shape) x_arr[setitem_indices] = rhs_arr x[setitem_indices] = rhs_arr assert all_equal(x, x_arr) # Using a list of lists rhs_list = (-np.ones(sliced_shape)).tolist() + x_arr = _as_numpy(x_arr) x_arr[setitem_indices] = rhs_list x[setitem_indices] = rhs_list assert all_equal(x, x_arr) @@ -1122,7 +1123,7 @@ def test_array_wrap_method(odl_tspace_impl): space = odl.tensor_space((3, 4), dtype='float32', exponent=1, weighting=2, impl=impl) x_arr, x = noise_elements(space) - y_arr = np.sin(x_arr) + y_arr = _module(tspace_impl).sin(x_arr) y = np.sin(x) # Should yield again an ODL tensor assert all_equal(y, y_arr) @@ -1732,11 +1733,12 @@ def test_ufunc_corner_cases(odl_tspace_impl): # Check that the result space is the same assert res.space == space - # Check usage of `order` argument + # Check usage of `order` argument (not available in cupy) for order in ('C', 'F'): - res = x.__array_ufunc__(np.sin, '__call__', x, order=order) - assert all_almost_equal(res, np.sin(x.asarray())) - assert res.data.flags[order + '_CONTIGUOUS'] + if tspace_impl == 'numpy': + res = x.__array_ufunc__(np.sin, '__call__', x, order=order) + assert all_almost_equal(res, np.sin(x.asarray())) + assert res.data.flags[order + '_CONTIGUOUS'] # Check usage of `dtype` argument res = x.__array_ufunc__(np.sin, '__call__', x, dtype='float32') @@ -1777,7 +1779,7 @@ def test_ufunc_corner_cases(odl_tspace_impl): res = x.__array_ufunc__(np.add, 'accumulate', x) assert all_almost_equal(res, np.add.accumulate(x.asarray())) assert res.space == space - arr = np.empty_like(x) + arr = _module(tspace_impl).empty_like(x) res = x.__array_ufunc__(np.add, 'accumulate', x, out=(arr,)) assert all_almost_equal(arr, np.add.accumulate(x.asarray())) assert res is arr @@ -1798,7 +1800,7 @@ def test_ufunc_corner_cases(odl_tspace_impl): assert all_almost_equal(res, np.add.reduce(x.asarray())) # With `out` argument and `axis` - out_ax0 = np.empty(3) + out_ax0 = _module(tspace_impl).empty(3) res = x.__array_ufunc__(np.add, 'reduce', x, axis=0, out=(out_ax0,)) assert all_almost_equal(out_ax0, np.add.reduce(x.asarray(), axis=0)) assert res is out_ax0 From 039f76f12afe003190ea03916c6605aaff721c25 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Tue, 28 Nov 2017 01:03:52 +0100 Subject: [PATCH 22/38] WIP: fix cupy stuff --- odl/space/cupy_tensors.py | 2 +- odl/test/space/tensors_test.py | 6 ++++++ 2 files changed, 7 insertions(+), 1 deletion(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index 60f9df5131b..0597589c3e9 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -2588,7 +2588,7 @@ def inner(self, x1, x2): if np.issubsctype(x1.dtype, np.complexfloating): dot = cupy.vdot(x2.data, x1.data) else: - dot = cupy.dot(x2.data, x1.data) + dot = cupy.dot(x2.data.ravel(), x1.data.ravel()) # complex(cupy_complex_scalar) not implemented, see # https://github.com/cupy/cupy/issues/782 diff --git a/odl/test/space/tensors_test.py b/odl/test/space/tensors_test.py index a6567cf3d34..da498c20be7 100644 --- a/odl/test/space/tensors_test.py +++ b/odl/test/space/tensors_test.py @@ -979,6 +979,9 @@ def test_transpose(odl_tspace_impl): assert x.T.is_linear # Check result + print(x.shape, x.dtype) + print(y.shape, y.dtype) + print('****************') assert x.T(y) == pytest.approx(y.inner(x)) assert all_equal(x.T.adjoint(1.0), x) @@ -1739,6 +1742,9 @@ def test_ufunc_corner_cases(odl_tspace_impl): res = x.__array_ufunc__(np.sin, '__call__', x, order=order) assert all_almost_equal(res, np.sin(x.asarray())) assert res.data.flags[order + '_CONTIGUOUS'] + elif tspace_impl == 'cupy': + with pytest.xfail(reason='cupy does not accept `order` in ufuncs'): + res = x.__array_ufunc__(np.sin, '__call__', x, order=order) # Check usage of `dtype` argument res = x.__array_ufunc__(np.sin, '__call__', x, dtype='float32') From 07822cc482fdb3573badf805a79ce43fb5969fc5 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Tue, 28 Nov 2017 19:00:19 +0100 Subject: [PATCH 23/38] WIP: fix cupy tensor tests --- odl/space/cupy_tensors.py | 294 ++++++++++++++++++++------------- odl/test/space/tensors_test.py | 153 +++++++++-------- odl/util/ufuncs.py | 21 ++- 3 files changed, 275 insertions(+), 193 deletions(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index 0597589c3e9..1f6f74624e3 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -1802,30 +1802,40 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): else: out1 = out_tuple[0] if isinstance(out_tuple[1], type(self)): - out1 = out_tuple[1].data + out2 = out_tuple[1].data else: - out1 = out_tuple[1] + out2 = out_tuple[1] # --- Process `inputs` --- # # Pull out the data container of the inputs if necessary - inputs = tuple( - inp.data if isinstance(inp, type(self)) else inp - for inp in inputs) + inputs = [inp.data if isinstance(inp, type(self)) else inp + for inp in inputs] # Determine native ufunc vs. Numpy ufunc if any(isinstance(o, np.ndarray) for o in out_tuple): native_ufunc = None use_native = False - elif any(not isinstance(inp, type(self.data)) for inp in inputs): + elif not all(isinstance(inp, type(self.data)) or np.isscalar(inp) + for inp in inputs): native_ufunc = None use_native = False else: native_ufunc = getattr(cupy, ufunc.__name__, None) - # Manual assignment for sum, cumsum, prod and cumprod - if (ufunc in (np.add, np.multiply) and - method in ('reduce', 'accumulate')): - use_native = native_ufunc is not None + + # Manually implemented cases + if ( + (ufunc in (np.add, np.multiply) and + method in ('reduce', 'accumulate') + ) or + (ufunc in (np.minimum, np.maximum) and + method == 'reduce' + ) or + (ufunc == np.add and + method == 'at' + ) + ): + use_native = native_ufunc is not None # should be True else: use_native = (native_ufunc is not None and hasattr(native_ufunc, method)) @@ -1843,35 +1853,63 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): # since cupy ufuncs do not accept array-like input. See # https://github.com/cupy/cupy/issues/594 and # https://github.com/odlgroup/odl/issues/1248 - inputs, orig_inputs = [], inputs - for inp in orig_inputs: - if (isinstance(inp, cupy.ndarray) or - np.isscalar(inp) or - inp is None): - inputs.append(inp) - else: - inputs.append(cupy.array(inp)) + for i in range(len(inputs)): + if not (isinstance(inputs[i], cupy.ndarray) or + np.isscalar(inputs[i]) or + inputs[i] is None): + inputs[i] = cupy.array(inputs[i]) + assert native_ufunc is not None + elif not use_native and method != 'at': # TODO: remove when upstream issue is fixed - # For non-native ufuncs (except `at`), we need ot cast our tensors + # For non-native ufuncs, we need ot cast our tensors # and Cupy arrays to Numpy arrays explicitly, since `__array__` # and friends are not implemented. See # https://github.com/cupy/cupy/issues/589 and # https://github.com/odlgroup/odl/issues/1248 - inputs, orig_inputs = [], inputs - for inp in orig_inputs: - if isinstance(inp, cupy.ndarray): - inputs.append(cupy.asnumpy(inp)) - elif isinstance(inp, CupyTensor): - inputs.append(cupy.asnumpy(inp.data)) + # We must exclude 'at' since it's in-place, thus we're not allowed + # to change the input. + for i in range(len(inputs)): + if isinstance(inputs[i], cupy.ndarray): + inputs[i] = cupy.asnumpy(inputs[i]) + elif isinstance(inputs[i], CupyTensor): + inputs[i] = cupy.asnumpy(inputs[i].data) + + # For debugging + if use_native: + assert all(isinstance(i, cupy.ndarray) or + np.isscalar(i) or + i is None + for i in inputs) + assert all(isinstance(o, cupy.ndarray) or o is None + for o in (out1, out2)) + + elif not use_native and method != 'at': + assert not any(isinstance(i, cupy.ndarray) for i in inputs) + # `out` handled below, we don't want to mess with it here + + # --- For later --- # + + # Wrap the result in an appropriate space, propagating weighting + # if possible + def space_element(res, use_weighting=True): + if use_weighting and is_floating_dtype(res.dtype): + if res.shape == self.shape: + weighting = self.space.weighting else: - inputs.append(inp) + # Don't propagate weighting if shape changes + # (but keep the exponent) + weighting = CupyTensorSpaceConstWeighting( + 1.0, self.space.exponent) - # --- Get some parameters for later --- # + spc_kwargs = {'weighting': weighting} + else: + # No `exponent` or `weighting` applicable + spc_kwargs = {} - # Arguments for space constructors - exponent = self.space.exponent - weighting = self.space.weighting + res_space = type(self.space)( + res.shape, res.dtype, self.device, **spc_kwargs) + return res_space.element(res) # --- Evaluate ufunc --- # @@ -1881,51 +1919,67 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): kwargs['out'] = out # No tuple packing for cupy res = native_ufunc(*inputs, **kwargs) else: - kwargs['out'] = (out,) - # Everything is cast to Numpy arrays by the parent method; - # the result can be a Numpy array or a tensor + # We need to explicitly cast `out` to Numpy array since + # the upstream `np.asarray` won't work. See + # https://github.com/cupy/cupy/issues/589 and + # https://github.com/odlgroup/odl/issues/1248 + if isinstance(out, cupy.ndarray): + out_npy = cupy.asnumpy(out) + else: + out_npy = out + + kwargs['out'] = (out_npy,) res = super(CupyTensor, self).__array_ufunc__( ufunc, '__call__', *inputs, **kwargs) - # Wrap result if necessary (lazily) + if isinstance(out, cupy.ndarray): + out[:] = cupy.asarray(res) + if out is None: - if is_floating_dtype(res.dtype): - # Weighting contains exponent - spc_kwargs = {'weighting': weighting} - else: - # No `exponent` or `weighting` applicable - spc_kwargs = {} - out_space = type(self.space)( - self.shape, res.dtype, self.device, **spc_kwargs) - return out_space.element(res) + result = space_element(res, use_weighting=True) else: - # `out` may be the unwrapped version, return the original - return out_tuple[0] + result = out_tuple[0] + + return result elif ufunc.nout == 2: - kwargs['out'] = (out1, out2) if use_native: - res1, res2 = native_ufunc(*inputs, **kwargs) + # cupy ufuncs with 2 outputs don't support writing to + # `out` + res1, res2 = native_ufunc(*inputs) + if out1 is not None: + out1[:] = res1 + if out2 is not None: + out2[:] = res2 else: - # Everything is cast to Numpy arrays by the parent method; - # the results can be Numpy arrays or tensors + # See case nout == 1 for comments + if isinstance(out1, cupy.ndarray): + out1_npy = cupy.asnumpy(out1) + else: + out1_npy = out1 + if isinstance(out2, cupy.ndarray): + out2_npy = cupy.asnumpy(out2) + else: + out2_npy = out2 + + kwargs['out'] = (out1_npy, out2_npy) res1, res2 = super(CupyTensor, self).__array_ufunc__( ufunc, '__call__', *inputs, **kwargs) - # Wrap results if necessary (lazily) - # We don't use exponents or weightings since we don't know + if isinstance(out1, cupy.ndarray): + out1[:] = cupy.asarray(res1) + if isinstance(out2, cupy.ndarray): + out2[:] = cupy.asarray(res2) + + # Don't use exponents or weightings since we don't know # how to map them to the spaces if out1 is None: - res_space = type(self.space)( - self.shape, res1.dtype, self.device) - result1 = res_space.element(res1) + result1 = space_element(res1, use_weighting=False) else: result1 = out_tuple[0] if out2 is None: - res_space = type(self.space)( - self.shape, res2.dtype, self.device) - result2 = res_space.element(res2) + result2 = space_element(res2, use_weighting=False) else: result2 = out_tuple[1] @@ -1935,6 +1989,53 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): raise NotImplementedError('nout = {} not supported' ''.format(ufunc.nout)) + # Special case 1 + elif (use_native and + (ufunc in (np.add, np.multiply) and + method in ('reduce', 'accumulate') + ) or + (ufunc in (np.minimum, np.maximum) and + method == 'reduce') + ): + # These cases are implemented but not available as methods + # of cupy.ufunc. We map the implementation by hand. + if ufunc == np.add: + function = cupy.sum if method == 'reduce' else cupy.cumsum + elif ufunc == np.multiply: + function = cupy.prod if method == 'reduce' else cupy.cumprod + elif ufunc == np.minimum and method == 'reduce': + function = cupy.min + elif ufunc == np.maximum and method == 'reduce': + function = cupy.max + else: + raise RuntimeError('no native {}.{}' + ''.format(ufunc.__name__, method)) + + # Make 0 the default for `axis` as the ufunc methods. The default + # is to reduce/accumulate over all axes. + axis = kwargs.pop('axis', 0) + kwargs['axis'] = axis + kwargs['out'] = out + res = function(*inputs, **kwargs) + + # Shortcut for scalar return value + if getattr(res, 'shape', ()) == (): + # Happens for `reduce` with all axes + return _python_scalar(res) + + if out is None: + result = space_element(res, use_weighting=True) + else: + result = out_tuple[0] + + return result + + # Special case 2 + elif use_native and ufunc == np.add and method == 'at': + cupy.scatter_add(*inputs, **kwargs) + return + + # Separate handling of 'at' since input is unchanged elif method == 'at': native_method = getattr(native_ufunc, 'at', None) use_native = (use_native and native_method is not None) @@ -1979,45 +2080,7 @@ def eval_at_via_npy(*inputs, **kwargs): else: eval_at_via_npy(*inputs, **kwargs) - elif (ufunc in (np.add, np.multiply) and - method in ('reduce', 'accumulate')): - # These cases are implemented but not available as methods - # of cupy.ufunc. We map the implementation by hand. - if ufunc == np.add: - function = cupy.sum if method == 'reduce' else cupy.cumsum - else: - function = cupy.prod if method == 'reduce' else cupy.cumprod - - # Make 0 the default for `axis` as the ufunc methods. The default - # for `[cum]sum` and `[cum]prod` is to reduce over all axes - axis = kwargs.pop('axis', 0) - kwargs['axis'] = axis - kwargs['out'] = out - res = function(*inputs, **kwargs) - - # Shortcut for scalar return value - if getattr(res, 'shape', ()) == (): - # Happens for `reduce` with all axes - return _python_scalar(res) - - # Wrap result if necessary (lazily) - if out is None: - if is_floating_dtype(res.dtype): - if res.shape != self.shape: - # Don't propagate weighting if shape changes - weighting = CupyTensorSpaceConstWeighting( - 1.0, exponent) - spc_kwargs = {'weighting': weighting} - else: - spc_kwargs = {} - - res_space = type(self.space)( - res.shape, res.dtype, self.device, **spc_kwargs) - result = res_space.element(res) - else: - result = out_tuple[0] - - return result + return else: # method != '__call__' kwargs['out'] = (out,) @@ -2025,7 +2088,7 @@ def eval_at_via_npy(*inputs, **kwargs): use_native = (use_native and native_method is not None) if use_native: - # Native method could exist but raise `NotImplementedError` + # A native method could exist but raise `NotImplementedError` # or return `NotImplemented`. We fall back to Numpy also in # that situation. try: @@ -2039,29 +2102,26 @@ def eval_at_via_npy(*inputs, **kwargs): ufunc, method, *inputs, **kwargs) else: + # See case `method == '__call__', nout == 1` for details + if isinstance(out, cupy.ndarray): + out_npy = cupy.asnumpy(out) + else: + out_npy = out + + kwargs['out'] = (out_npy,) res = super(CupyTensor, self).__array_ufunc__( ufunc, method, *inputs, **kwargs) + if isinstance(out, cupy.ndarray): + out[:] = cupy.asarray(res) + # Shortcut for scalar or no return value - if np.isscalar(res) or res is None: - # The first occurs for `reduce` with all axes, - # the second for in-place stuff (`at` currently) - return res + if getattr(res, 'shape', ()) == (): + # Occurs for `reduce` with all axes + return _python_scalar(res) - # Wrap result if necessary (lazily) if out is None: - if is_floating_dtype(res.dtype): - if res.shape != self.shape: - # Don't propagate weighting if shape changes - weighting = CupyTensorSpaceConstWeighting( - 1.0, exponent) - spc_kwargs = {'weighting': weighting} - else: - spc_kwargs = {} - - res_space = type(self.space)( - res.shape, res.dtype, self.device, **spc_kwargs) - result = res_space.element(res) + result = space_element(res, use_weighting=True) else: result = out_tuple[0] @@ -2190,7 +2250,7 @@ def conj(self, out=None): return self.space.element(self.data.conj()) else: if self.space.is_real: - self.assign(out) + out.assign(self) else: # In-place not available as it seems out[:] = self.data.conj() @@ -2200,7 +2260,7 @@ def __ipow__(self, other): """Return ``self **= other``.""" try: if other == int(other): - return super(CupyTensorSpace, self).__ipow__(other) + return super(CupyTensor, self).__ipow__(other) except TypeError: pass @@ -2588,6 +2648,8 @@ def inner(self, x1, x2): if np.issubsctype(x1.dtype, np.complexfloating): dot = cupy.vdot(x2.data, x1.data) else: + # Ravels in C order, no other supported currently + # TODO: ravel in most efficient ordering when available dot = cupy.dot(x2.data.ravel(), x1.data.ravel()) # complex(cupy_complex_scalar) not implemented, see diff --git a/odl/test/space/tensors_test.py b/odl/test/space/tensors_test.py index da498c20be7..53ba20441a7 100644 --- a/odl/test/space/tensors_test.py +++ b/odl/test/space/tensors_test.py @@ -17,13 +17,14 @@ import odl from odl.set.space import LinearSpaceTypeError +from odl.space.entry_points import tensor_space_impl from odl.space.npy_tensors import ( - NumpyTensor, NumpyTensorSpace, + NumpyTensorSpace, NumpyTensorSpaceConstWeighting, NumpyTensorSpaceArrayWeighting, NumpyTensorSpaceCustomInner, NumpyTensorSpaceCustomNorm, NumpyTensorSpaceCustomDist) from odl.space.cupy_tensors import ( - CupyTensor, CupyTensorSpace, + CupyTensorSpace, CupyTensorSpaceConstWeighting, CupyTensorSpaceArrayWeighting, CupyTensorSpaceCustomInner, CupyTensorSpaceCustomNorm, CupyTensorSpaceCustomDist, @@ -84,16 +85,6 @@ def _pos_array(space): return _module(space.impl).abs(noise_array(space)) + 0.1 -def _odl_tensor_cls(impl): - """Return the ODL tensor class for given impl.""" - if impl == 'numpy': - return NumpyTensor - elif impl == 'cupy': - return CupyTensor - else: - assert False - - def _weighting_cls(impl, kind): """Return the weighting class for given impl and kind.""" if impl == 'numpy': @@ -152,8 +143,8 @@ def weight(request): @pytest.fixture(scope='module') def tspace(floating_dtype, odl_tspace_impl): - cls = odl.space.entry_points.tensor_space_impl(odl_tspace_impl) - if floating_dtype not in cls.available_dtypes(): + available_dtypes = tensor_space_impl(odl_tspace_impl).available_dtypes() + if floating_dtype not in available_dtypes: pytest.skip('dtype {} not supported by impl {!r}' ''.format(floating_dtype, odl_tspace_impl)) else: @@ -863,7 +854,7 @@ def test_element_getitem(odl_tspace_impl, getitem_indices): sliced_shape = x_arr_sliced.shape x_sliced = x[getitem_indices] - if np.isscalar(x_arr_sliced): + if np.isscalar(x_sliced): assert x_arr_sliced == x_sliced else: assert x_sliced.shape == sliced_shape @@ -1525,38 +1516,16 @@ def _check_result_is_out(result, out_seq): else: assert False - if (np.issubsctype(tspace.dtype, np.floating) or - np.issubsctype(tspace.dtype, np.complexfloating) and - name in ['bitwise_and', - 'bitwise_or', - 'bitwise_xor', - 'invert', - 'left_shift', - 'right_shift']): - # Skip integer only methods for floating point data types - return - - if (np.issubsctype(tspace.dtype, np.complexfloating) and - name in ['remainder', - 'trunc', - 'signbit', - 'invert', - 'left_shift', - 'right_shift', - 'rad2deg', - 'deg2rad', - 'copysign', - 'mod', - 'modf', - 'fmod', - 'logaddexp2', - 'logaddexp', - 'hypot', - 'arctan2', - 'floor', - 'ceil']): - # Skip real-only methods for complex data types - return + # See https://github.com/cupy/cupy/issues/794 + cupy_ufuncs_broken_complex = [ + 'expm1', 'floor_divide', 'fmin', 'fmax', + 'greater', 'greater_equal', 'less', 'less_equal', 'log1p', 'log2', + 'logical_and', 'logical_or', 'logical_not', 'logical_xor', 'minimum', + 'maximum', 'rint', 'sign', 'square'] + if (tspace.impl == 'cupy' and + tspace.dtype.kind == 'c' and + ufunc in cupy_ufuncs_broken_complex): + pytest.xfail('ufunc {} broken for complex input in cupy'.format(ufunc)) # Create some data arrays, elements = noise_elements(tspace, nin + nout) @@ -1572,20 +1541,30 @@ def _check_result_is_out(result, out_seq): out_arr_kwargs = {'out': out_arrays[:nout]} out_elem_kwargs = {'out': out_elems[:nout]} - # Get function to call, using both interfaces: # - vec.ufunc(*other_args) # - np.ufunc(vec, *other_args) ufunc_method = getattr(data_elem.ufuncs, name) in_elems_method = elements[1:nin] in_elems_npy = elements[:nin] - result_npy = ufunc_npy(*in_arrays_npy) + try: + result_npy = ufunc_npy(*in_arrays_npy) + except TypeError: + pytest.xfail('numpy ufunc not valid for inputs') # Out-of-place -- in = elements -- ufunc = method or numpy result = ufunc_method(*in_elems_method) assert all_almost_equal(result_npy, result) _check_result_type(result, tspace.element_type) + # Get element(s) in the right space for in-place later + if nout == 1: + out_elems = [result.space.element()] + elif nout == 2: + out_elems = [res.space.element() for res in result] + else: + assert False + result = ufunc_npy(*in_elems_npy) assert all_almost_equal(result_npy, result) _check_result_type(result, tspace.element_type) @@ -1603,9 +1582,18 @@ def _check_result_is_out(result, out_seq): result = ufunc_method(*in_elems_method, **out_elem_kwargs) assert all_almost_equal(result_npy, result) _check_result_is_out(result, out_elems[:nout]) + if USE_ARRAY_UFUNCS_INTERFACE: # Custom objects not allowed as `out` for numpy < 1.13 result = ufunc_npy(*in_elems_npy, **out_elem_kwargs) + + if nout == 1: + kwargs_out = {'out': out_elems[0]} + elif nout == 2: + kwargs_out = {'out': (out_elems[0], out_elems[1])} + + result = ufunc_npy(*in_elems_npy, **kwargs_out) + assert all_almost_equal(result_npy, result) _check_result_is_out(result, out_elems[:nout]) @@ -1624,7 +1612,11 @@ def _check_result_is_out(result, out_seq): kwargs_npy = {'out': out_arrays_npy[:nout]} kwargs_own = {'out': out_arrays_own[:nout]} - result_out_npy = ufunc_npy(*in_elems_npy, **kwargs_npy) + try: + result_out_npy = ufunc_npy(*in_elems_npy, **kwargs_npy) + except TypeError: + pytest.xfail('numpy ufunc not valid for inputs') + result_out_own = ufunc_npy(*in_elems_npy, **kwargs_own) assert all_almost_equal(result_out_npy, result_npy) assert all_almost_equal(result_out_own, result_npy) @@ -1639,12 +1631,21 @@ def _check_result_is_out(result, out_seq): mod_array = in_arrays_npy[0].copy() mod_elem = in_elems_npy[0].copy() if nin == 1: - result_npy = ufunc_npy.at(mod_array, indices) + try: + result_npy = ufunc_npy.at(mod_array, indices) + except TypeError: + pytest.xfail('numpy ufunc.at not valid for inputs') + result = ufunc_npy.at(mod_elem, indices) + elif nin == 2: other_array = in_arrays_npy[1][indices] other_elem = in_elems_npy[1][indices] - result_npy = ufunc_npy.at(mod_array, indices, other_array) + try: + result_npy = ufunc_npy.at(mod_array, indices, other_array) + except TypeError: + pytest.xfail('numpy ufunc.at not valid for inputs') + result = ufunc_npy.at(mod_elem, indices, other_elem) assert all_almost_equal(result, result_npy) @@ -1657,8 +1658,12 @@ def _check_result_is_out(result, out_seq): # We only test along one axis since some binary ufuncs are not # re-orderable, in which case Numpy raises a ValueError + try: + result_npy = ufunc_npy.reduce(in_array) + except TypeError: + pytest.xfail('numpy ufunc.reduce not valid for inputs') + # Out-of-place -- in = element - result_npy = ufunc_npy.reduce(in_array) result = ufunc_npy.reduce(in_elem) assert all_almost_equal(result, result_npy) result_keepdims = ufunc_npy.reduce(in_elem, keepdims=True) @@ -1678,7 +1683,14 @@ def _check_result_is_out(result, out_seq): assert all_almost_equal(out_array_own, result_keepdims) # Using a specific dtype - result_npy = ufunc_npy.reduce(in_array, dtype=complex) + try: + result_npy = ufunc_npy.reduce(in_array, dtype=complex) + except TypeError: + pytest.xfail('numpy ufunc.reduce not valid for complex dtype') + + if tspace.impl == 'cupy': + pytest.xfail('cupy ufunc.reduce raises error for complex dtype') + result = ufunc_npy.reduce(in_elem, dtype=complex) assert result.dtype == result_npy.dtype assert all_almost_equal(result, result_npy) @@ -1846,51 +1858,60 @@ def testodl_reduction(tspace, odl_reduction): # Should be equal theoretically, but summation order, other stuff, ..., # hence we use approx + if (tspace.impl == 'cupy' and + reduction in ('min', 'max') and + tspace.dtype.kind == 'c'): + pytest.xfail('Cupy does not accept complex input to `min` and `max`') + # Full reduction, produces scalar - result_npy = npy_reduction(x_arr) + result_npy = npy_reduction(_as_numpy(x_arr)) result = x_reduction() assert result == pytest.approx(result_npy) result = x_reduction(axis=(0, 1)) assert result == pytest.approx(result_npy) # Reduction along axes, produces element in reduced space - result_npy = npy_reduction(x_arr, axis=0) + result_npy = npy_reduction(_as_numpy(x_arr), axis=0) result = x_reduction(axis=0) - assert isinstance(result, NumpyTensor) + assert isinstance(result, tspace.element_type) assert result.shape == result_npy.shape assert result.dtype == x.dtype - assert np.allclose(result, result_npy) + assert all_almost_equal(result, result_npy) # Check reduced space properties - assert isinstance(result.space, NumpyTensorSpace) + assert type(result.space) is type(tspace) assert result.space.exponent == x.space.exponent assert result.space.weighting == x.space.weighting # holds true here # Evaluate in-place out = result.space.element() x_reduction(axis=0, out=out) - assert np.allclose(out, result_npy) + assert all_almost_equal(out, result_npy) # Use keepdims parameter - result_npy = npy_reduction(x_arr, axis=1, keepdims=True) + result_npy = npy_reduction(_as_numpy(x_arr), axis=1, keepdims=True) result = x_reduction(axis=1, keepdims=True) assert result.shape == result_npy.shape - assert np.allclose(result, result_npy) + assert all_almost_equal(result, result_npy) # Evaluate in-place out = result.space.element() x_reduction(axis=1, keepdims=True, out=out) - assert np.allclose(out, result_npy) + assert all_almost_equal(out, result_npy) # Use dtype parameter # These reductions have a `dtype` parameter if name in ('cumprod', 'cumsum', 'mean', 'prod', 'std', 'sum', 'trace', 'var'): - result_npy = npy_reduction(x_arr, axis=1, dtype='complex64') + if tspace.impl == 'cupy' and tspace.dtype == 'float16': + # See https://github.com/cupy/cupy/issues/795 + pytest.xfail('reduction with complex fails for float16 in cupy') + + result_npy = npy_reduction(_as_numpy(x_arr), axis=1, dtype='complex64') result = x_reduction(axis=1, dtype='complex64') assert result.dtype == np.dtype('complex64') - assert np.allclose(result, result_npy) + assert all_almost_equal(result, result_npy) # Evaluate in-place out = result.space.element() x_reduction(axis=1, dtype='complex64', out=out) - assert np.allclose(out, result_npy) + assert all_almost_equal(out, result_npy) def test_ufunc_reduction_docs_notempty(odl_tspace_impl): diff --git a/odl/util/ufuncs.py b/odl/util/ufuncs.py index d2694553140..d4d5d80dd5d 100644 --- a/odl/util/ufuncs.py +++ b/odl/util/ufuncs.py @@ -96,12 +96,9 @@ def wrapper(self, out=None, **kwargs): ufunc, '__call__', self.elem, out=out, **kwargs) elif n_out == 2: - def wrapper(self, out=None, **kwargs): - if out is None: - out = (None, None) - + def wrapper(self, out1=None, out2=None, **kwargs): return self.elem.__array_ufunc__( - ufunc, '__call__', self.elem, out=out, **kwargs) + ufunc, '__call__', self.elem, out=(out1, out2), **kwargs) else: raise NotImplementedError @@ -109,16 +106,18 @@ def wrapper(self, out=None, **kwargs): elif n_in == 2: if n_out == 1: def wrapper(self, x2, out=None, **kwargs): + if out is None or isinstance(out, (type(self.elem), + type(self.elem.data))): + out = (out,) + return self.elem.__array_ufunc__( - ufunc, '__call__', self.elem, x2, out=(out,), **kwargs) + ufunc, '__call__', self.elem, x2, out=out, **kwargs) elif n_out == 2: - def wrapper(self, x2, out=None, **kwargs): - if out is None: - out = (None, None) - + def wrapper(self, x2, out1=None, out2=None, **kwargs): return self.elem.__array_ufunc__( - ufunc, '__call__', self.elem, x2, out=out, **kwargs) + ufunc, '__call__', self.elem, x2, out=(out1, out2), + **kwargs) else: raise NotImplementedError From 47f74c98d65f56dc54fdbdf9fec3ce9da2c74782 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Wed, 29 Nov 2017 08:09:14 +0100 Subject: [PATCH 24/38] WIP: fix pspace noise_array and out aliasing issue --- odl/test/space/pspace_test.py | 5 +- odl/test/space/tensors_test.py | 105 ++++++++++++++++++--------------- odl/util/testutils.py | 25 ++++++-- odl/util/ufuncs.py | 98 +++++++++++++++++++++++------- 4 files changed, 159 insertions(+), 74 deletions(-) diff --git a/odl/test/space/pspace_test.py b/odl/test/space/pspace_test.py index ad7acd1ad51..a28e310ebf0 100644 --- a/odl/test/space/pspace_test.py +++ b/odl/test/space/pspace_test.py @@ -160,7 +160,7 @@ def test_is_power_space(): def test_mixed_space(): - """Verify that a mixed productspace is handled properly.""" + """Verify that a mixed product space is handled properly.""" r2_1 = odl.rn(2, dtype='float64') r2_2 = odl.rn(2, dtype='float32') pspace = odl.ProductSpace(r2_1, r2_2) @@ -837,8 +837,7 @@ def test_element_setitem_broadcast(): def test_unary_ops(): - # Verify that the unary operators (`+x` and `-x`) work as expected - + """Verify that the unary operators ``+x`` and ``-x`` work as expected.""" space = odl.rn(3) pspace = odl.ProductSpace(space, 2) diff --git a/odl/test/space/tensors_test.py b/odl/test/space/tensors_test.py index 53ba20441a7..d1d1347b328 100644 --- a/odl/test/space/tensors_test.py +++ b/odl/test/space/tensors_test.py @@ -279,8 +279,6 @@ def test_init_tspace_weighting(weight, exponent, odl_tspace_impl): weighting = weighting_cls(weight, exponent) - assert space.weighting == weighting - # Using a weighting instance space = odl.tensor_space((3, 4), weighting=weighting, exponent=exponent, impl=impl) @@ -607,7 +605,7 @@ def test_multiply(tspace): def test_multiply_exceptions(tspace): """Test if multiply raises correctly for bad input.""" - other_space = odl.rn((4, 3)) + other_space = odl.rn((4, 3), impl=tspace.impl) other_x = other_space.zero() x, y = tspace.zero(), tspace.zero() @@ -748,16 +746,16 @@ def test_inner(tspace): """Test the inner method against numpy.vdot.""" xd = noise_element(tspace) yd = noise_element(tspace) - - # TODO: add weighting correct_inner = np.vdot(yd, xd) - assert tspace.inner(xd, yd) == pytest.approx(correct_inner) - assert xd.inner(yd) == pytest.approx(correct_inner) + + # Allow some error for single and half precision + assert tspace.inner(xd, yd) == pytest.approx(correct_inner, rel=1e-2) + assert xd.inner(yd) == pytest.approx(correct_inner, rel=1e-2) def test_inner_exceptions(tspace): """Test if inner raises correctly for bad input.""" - other_space = odl.rn((4, 3)) + other_space = odl.rn((4, 3), impl=tspace.impl) other_x = other_space.zero() x = tspace.zero() @@ -771,15 +769,16 @@ def test_inner_exceptions(tspace): def test_norm(tspace): """Test the norm method against numpy.linalg.norm.""" xarr, x = noise_elements(tspace) - correct_norm = np.linalg.norm(_as_numpy(xarr.ravel())) - assert tspace.norm(x) == pytest.approx(correct_norm) - assert x.norm() == pytest.approx(correct_norm) + + # Allow some error for single and half precision + assert tspace.norm(x) == pytest.approx(correct_norm, rel=1e-2) + assert x.norm() == pytest.approx(correct_norm, rel=1e-2) def test_norm_exceptions(tspace): """Test if norm raises correctly for bad input.""" - other_space = odl.rn((4, 3)) + other_space = odl.rn((4, 3), impl=tspace.impl) other_x = other_space.zero() with pytest.raises(LinearSpaceTypeError): @@ -806,15 +805,16 @@ def test_pnorm(exponent, odl_tspace_impl): def test_dist(tspace): """Test the dist method against numpy.linalg.norm of the difference.""" [xarr, yarr], [x, y] = noise_elements(tspace, n=2) - correct_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel())) - assert tspace.dist(x, y) == pytest.approx(correct_dist) + + # Allow some error for single and half precision + assert tspace.dist(x, y) == pytest.approx(correct_dist, rel=1e-2) assert x.dist(y) == pytest.approx(correct_dist) def test_dist_exceptions(tspace): """Test if dist raises correctly for bad input.""" - other_space = odl.rn((4, 3)) + other_space = odl.rn((4, 3), impl=tspace.impl) other_x = other_space.zero() x = tspace.zero() @@ -889,12 +889,12 @@ def test_element_setitem(odl_tspace_impl, setitem_indices): assert all_equal(x, x_arr) # Setting values with arrays - rhs_arr = _module(tspace_impl).ones(sliced_shape) + rhs_arr = _module(impl).ones(sliced_shape) x_arr[setitem_indices] = rhs_arr x[setitem_indices] = rhs_arr assert all_equal(x, x_arr) - # Using a list of lists + # Setting values with a list of lists rhs_list = (-np.ones(sliced_shape)).tolist() x_arr = _as_numpy(x_arr) x_arr[setitem_indices] = rhs_list @@ -970,9 +970,6 @@ def test_transpose(odl_tspace_impl): assert x.T.is_linear # Check result - print(x.shape, x.dtype) - print(y.shape, y.dtype) - print('****************') assert x.T(y) == pytest.approx(y.inner(x)) assert all_equal(x.T.adjoint(1.0), x) @@ -1036,7 +1033,9 @@ def test_conversion_to_scalar(odl_tspace_impl): """Test conversion of size-1 vectors/tensors to scalars.""" impl = odl_tspace_impl space = odl.rn(1, impl=impl) + # Size 1 real space + space = odl.rn(1, impl=impl) value = 1.5 element = space.element(value) @@ -1117,7 +1116,7 @@ def test_array_wrap_method(odl_tspace_impl): space = odl.tensor_space((3, 4), dtype='float32', exponent=1, weighting=2, impl=impl) x_arr, x = noise_elements(space) - y_arr = _module(tspace_impl).sin(x_arr) + y_arr = _module(impl).sin(x_arr) y = np.sin(x) # Should yield again an ODL tensor assert all_equal(y, y_arr) @@ -1137,7 +1136,7 @@ def test_conj(tspace): assert all_equal(y, xarr.conj()) -# --- Weightings (Numpy) --- # +# --- Weightings --- # def test_array_weighting_init(odl_tspace_impl, exponent): @@ -1241,8 +1240,9 @@ def test_array_weighting_inner(tspace): weight_arr = _pos_array(tspace) weighting_cls = _weighting_cls(tspace.impl, 'array') weighting = weighting_cls(weight_arr) - true_inner = np.vdot(_as_numpy(yarr), _as_numpy(xarr * weight_arr)) + + # Allow some error for single and half precision assert weighting.inner(x, y) == pytest.approx(true_inner, rel=1e-2) # Exponent != 2 -> no inner product, should raise @@ -1257,7 +1257,6 @@ def test_array_weighting_norm(tspace, exponent): weight_arr = _pos_array(tspace) weighting_cls = _weighting_cls(tspace.impl, 'array') weighting = weighting_cls(weight_arr, exponent=exponent) - if exponent == float('inf'): true_norm = np.linalg.norm(_as_numpy(xarr.ravel()), ord=float('inf')) else: @@ -1265,6 +1264,7 @@ def test_array_weighting_norm(tspace, exponent): _as_numpy((weight_arr ** (1 / exponent) * xarr).ravel()), ord=exponent) + # Allow some error for single and half precision assert weighting.norm(x) == pytest.approx(true_norm, rel=1e-2) @@ -1275,7 +1275,6 @@ def test_array_weighting_dist(tspace, exponent): weight_arr = _pos_array(tspace) weighting_cls = _weighting_cls(tspace.impl, 'array') weighting = weighting_cls(weight_arr, exponent=exponent) - if exponent == float('inf'): true_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel()), ord=float('inf')) @@ -1284,6 +1283,7 @@ def test_array_weighting_dist(tspace, exponent): _as_numpy((weight_arr ** (1 / exponent) * (xarr - yarr)).ravel()), ord=exponent) + # Allow some error for single and half precision assert weighting.dist(x, y) == pytest.approx(true_dist, rel=1e-2) @@ -1350,8 +1350,9 @@ def test_const_weighting_inner(tspace): constant = 1.5 weighting_cls = _weighting_cls(tspace.impl, 'const') weighting = weighting_cls(constant) - true_inner = constant * np.vdot(_as_numpy(yarr), _as_numpy(xarr)) + + # Allow some error for single and half precision assert weighting.inner(x, y) == pytest.approx(true_inner, rel=1e-2) # Exponent != 2 -> no inner @@ -1367,12 +1368,13 @@ def test_const_weighting_norm(tspace, exponent): constant = 1.5 weighting_cls = _weighting_cls(tspace.impl, 'const') weighting = weighting_cls(constant, exponent=exponent) - if exponent == float('inf'): factor = 1.0 else: factor = constant ** (1 / exponent) true_norm = factor * np.linalg.norm(_as_numpy(xarr.ravel()), ord=exponent) + + # Allow some error for single and half precision assert weighting.norm(x) == pytest.approx(true_norm, rel=1e-2) @@ -1383,13 +1385,14 @@ def test_const_weighting_dist(tspace, exponent): constant = 1.5 weighting_cls = _weighting_cls(tspace.impl, 'const') weighting = weighting_cls(constant, exponent=exponent) - if exponent == float('inf'): factor = 1.0 else: factor = constant ** (1 / exponent) true_dist = factor * np.linalg.norm(_as_numpy((xarr - yarr).ravel()), ord=exponent) + + # Allow some error for single and half precision assert weighting.dist(x, y) == pytest.approx(true_dist, rel=1e-2) @@ -1410,12 +1413,12 @@ def inner(x, y): assert w != w_other true_inner = np.vdot(_as_numpy(yarr), _as_numpy(xarr)) - assert w.inner(x, y) == pytest.approx(true_inner, rel=1e-2) - true_norm = np.linalg.norm(_as_numpy(xarr.ravel())) - assert w.norm(x) == pytest.approx(true_norm, rel=1e-2) - true_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel())) + + # Allow some error for single and half precision + assert w.inner(x, y) == pytest.approx(true_inner, rel=1e-2) + assert w.norm(x) == pytest.approx(true_norm, rel=1e-2) assert w.dist(x, y) == pytest.approx(true_dist, rel=1e-2) with pytest.raises(TypeError): @@ -1441,15 +1444,16 @@ def other_norm(x): assert w == w_same assert w != w_other - with pytest.raises(NotImplementedError): - w.inner(x, y) - true_norm = np.linalg.norm(_as_numpy(xarr.ravel())) - assert w.norm(x) == pytest.approx(true_norm, rel=1e-2) - true_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel())) + + # Allow some error for single and half precision + assert w.norm(x) == pytest.approx(true_norm, rel=1e-2) assert w.dist(x, y) == pytest.approx(true_dist, rel=1e-2) + with pytest.raises(NotImplementedError): + w.inner(x, y) + with pytest.raises(TypeError): weighting_cls(1) @@ -1473,15 +1477,17 @@ def other_dist(x, y): assert w == w_same assert w != w_other + true_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel())) + + # Allow some error for single and half precision + assert w.dist(x, y) == pytest.approx(true_dist, rel=1e-2) + with pytest.raises(NotImplementedError): w.inner(x, y) with pytest.raises(NotImplementedError): w.norm(x) - true_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel())) - assert w.dist(x, y) == pytest.approx(true_dist, rel=1e-2) - with pytest.raises(TypeError): weighting_cls(1) @@ -1498,6 +1504,10 @@ def testodl_ufuncs(tspace, odl_ufunc): nin = ufunc_npy.nin nout = ufunc_npy.nout + # Disable Numpy warnings for the time being + npy_err_orig = np.geterr() + np.seterr(all='ignore') + def _check_result_type(result, expected_type): if nout == 1: assert isinstance(result, tspace.element_type) @@ -1525,7 +1535,7 @@ def _check_result_is_out(result, out_seq): if (tspace.impl == 'cupy' and tspace.dtype.kind == 'c' and ufunc in cupy_ufuncs_broken_complex): - pytest.xfail('ufunc {} broken for complex input in cupy'.format(ufunc)) + pytest.xfail('ufunc {} broken for complex input in cupy'.format(name)) # Create some data arrays, elements = noise_elements(tspace, nin + nout) @@ -1698,6 +1708,9 @@ def _check_result_is_out(result, out_seq): # Other ufunc method use the same interface, to we don't perform # extra tests for them. + # Reset Numpy err handling + np.seterr(**npy_err_orig) + def test_ufunc_cupy_force_native(): """Test the ``force_native`` flag for cupy based ufuncs.""" @@ -1750,11 +1763,11 @@ def test_ufunc_corner_cases(odl_tspace_impl): # Check usage of `order` argument (not available in cupy) for order in ('C', 'F'): - if tspace_impl == 'numpy': + if impl == 'numpy': res = x.__array_ufunc__(np.sin, '__call__', x, order=order) assert all_almost_equal(res, np.sin(x.asarray())) assert res.data.flags[order + '_CONTIGUOUS'] - elif tspace_impl == 'cupy': + elif impl == 'cupy': with pytest.xfail(reason='cupy does not accept `order` in ufuncs'): res = x.__array_ufunc__(np.sin, '__call__', x, order=order) @@ -1797,7 +1810,7 @@ def test_ufunc_corner_cases(odl_tspace_impl): res = x.__array_ufunc__(np.add, 'accumulate', x) assert all_almost_equal(res, np.add.accumulate(x.asarray())) assert res.space == space - arr = _module(tspace_impl).empty_like(x) + arr = _module(impl).empty_like(x) res = x.__array_ufunc__(np.add, 'accumulate', x, out=(arr,)) assert all_almost_equal(arr, np.add.accumulate(x.asarray())) assert res is arr @@ -1818,7 +1831,7 @@ def test_ufunc_corner_cases(odl_tspace_impl): assert all_almost_equal(res, np.add.reduce(x.asarray())) # With `out` argument and `axis` - out_ax0 = _module(tspace_impl).empty(3) + out_ax0 = _module(impl).empty(3) res = x.__array_ufunc__(np.add, 'reduce', x, axis=0, out=(out_ax0,)) assert all_almost_equal(out_ax0, np.add.reduce(x.asarray(), axis=0)) assert res is out_ax0 @@ -1859,7 +1872,7 @@ def testodl_reduction(tspace, odl_reduction): # hence we use approx if (tspace.impl == 'cupy' and - reduction in ('min', 'max') and + name in ('min', 'max') and tspace.dtype.kind == 'c'): pytest.xfail('Cupy does not accept complex input to `min` and `max`') diff --git a/odl/util/testutils.py b/odl/util/testutils.py index deb4584ecf4..808aaa5b2e5 100644 --- a/odl/util/testutils.py +++ b/odl/util/testutils.py @@ -353,7 +353,23 @@ def noise_array(space): from odl.space.cupy_tensors import cupy if isinstance(space, ProductSpace): - return [noise_array(si) for si in space] + arr_list = [noise_array(si) for si in space] + try: + impl = space[0].impl + except (IndexError, AttributeError): + impl = 'numpy' + + if space.is_power_space: + if impl == 'numpy': + return np.vstack(arr_list) + elif impl == 'cupy': + return cupy.vstack(arr_list) + else: + raise RuntimeError('bad `impl` {!r}'.format(impl)) + + else: + return tuple(arr_list) + else: if space.dtype == bool: # TODO(kohr-h): use `randint(..., dtype=bool)` from Numpy 1.11 on @@ -371,12 +387,13 @@ def noise_array(space): raise ValueError('bad dtype {}'.format(space.dtype)) arr = arr.astype(space.dtype, copy=False) - if space.impl == 'numpy': + impl = getattr(space, 'impl', 'numpy') + if impl == 'numpy': return arr - elif space.impl == 'cupy': + elif impl == 'cupy': return cupy.asarray(arr) else: - raise RuntimeError('bad `impl` {!r}'.format(space.impl)) + raise RuntimeError('bad `impl` {!r}'.format(impl)) def noise_element(space): diff --git a/odl/util/ufuncs.py b/odl/util/ufuncs.py index d4d5d80dd5d..5e211f2a5a7 100644 --- a/odl/util/ufuncs.py +++ b/odl/util/ufuncs.py @@ -154,6 +154,18 @@ def sum(self, axis=None, dtype=None, out=None, keepdims=False): np.add, 'reduce', self.elem, axis=axis, dtype=dtype, out=(out,), keepdims=keepdims) + def cumsum(self, axis=None, dtype=None, out=None): + """Return the cumulative sum of ``self``. + + See Also + -------- + numpy.cumsum + cumprod + """ + return self.elem.__array_ufunc__( + np.add, 'accumulate', self.elem, + axis=axis, dtype=dtype, out=(out,)) + def prod(self, axis=None, dtype=None, out=None, keepdims=False): """Return the product of ``self``. @@ -166,6 +178,18 @@ def prod(self, axis=None, dtype=None, out=None, keepdims=False): np.multiply, 'reduce', self.elem, axis=axis, dtype=dtype, out=(out,), keepdims=keepdims) + def cumprod(self, axis=None, dtype=None, out=None): + """Return the cumulative product of ``self``. + + See Also + -------- + numpy.cumprod + cumsum + """ + return self.elem.__array_ufunc__( + np.multiply, 'accumulate', self.elem, + axis=axis, dtype=dtype, out=(out,)) + def min(self, axis=None, dtype=None, out=None, keepdims=False): """Return the minimum of ``self``. @@ -207,33 +231,49 @@ def wrap_ufunc_productspace(name, n_in, n_out, doc): def wrapper(self, out=None, **kwargs): from odl.space.pspace import ProductSpace if out is None: - out = [None] * len(self.elem.space) + out_seq = [None] * len(self.elem.space) + else: + out_seq = out res = [] - for xi, out_i in zip(self.elem, out): + for xi, out_i in zip(self.elem, out_seq): r = getattr(xi.ufuncs, name)(out=out_i, **kwargs) res.append(r) - out_space = ProductSpace(*[r.space for r in res]) - return out_space.element(res) + + if out is None: + out_space = ProductSpace(*[r.space for r in res]) + out = out_space.element(res) + + return out elif n_out == 2: def wrapper(self, out1=None, out2=None, **kwargs): from odl.space.pspace import ProductSpace if out1 is None: - out1 = [None] * len(self.elem.space) + out1_seq = [None] * len(self.elem.space) + else: + out1_seq = out1 if out2 is None: - out2 = [None] * len(self.elem.space) + out2_seq = [None] * len(self.elem.space) + else: + out2_seq = out2 res1, res2 = [], [] - for xi, out1_i, out2_i in zip(self.elem, out1, out2): + for xi, out1_i, out2_i in zip(self.elem, out1_seq, out2_seq): r1, r2 = getattr(xi.ufuncs, name)(out1=out1_i, out2=out2_i, **kwargs) res1.append(r1) res2.append(r2) - out_space_1 = ProductSpace(*[r.space for r in res1]) - out_space_2 = ProductSpace(*[r.space for r in res2]) - return out_space_1.element(res1), out_space_2.element(res2) + + if out1 is None: + out_space_1 = ProductSpace(*[r.space for r in res1]) + out1 = out_space_1.element(res1) + if out2 is None: + out_space_2 = ProductSpace(*[r.space for r in res2]) + out2 = out_space_2.element(res2) + + return out1, out2 else: raise NotImplementedError @@ -243,34 +283,50 @@ def wrapper(self, out1=None, out2=None, **kwargs): def wrapper(self, x2, out=None, **kwargs): from odl.space.pspace import ProductSpace if out is None: - out = [None] * len(self.elem.space) + out_seq = [None] * len(self.elem.space) + else: + out_seq = out res = [] - for x1_i, x2_i, out_i in zip(self.elem, x2, out): + for x1_i, x2_i, out_i in zip(self.elem, x2, out_seq): r = getattr(x1_i.ufuncs, name)(x2_i, out=out_i, **kwargs) res.append(r) - out_space = ProductSpace(*[r.space for r in res]) - return out_space.element(res) + + if out is None: + out_space = ProductSpace(*[r.space for r in res]) + out = out_space.element(res) + + return out elif n_out == 2: def wrapper(self, x2, out1=None, out2=None, **kwargs): from odl.space.pspace import ProductSpace if out1 is None: - out1 = [None] * len(self.elem.space) + out1_seq = [None] * len(self.elem.space) + else: + out1_seq = out1 if out2 is None: - out2 = [None] * len(self.elem.space) + out2_seq = [None] * len(self.elem.space) + else: + out2_seq = out2 res1, res2 = [], [] - for x1_i, x2_i, out1_i, out2_i in zip(self.elem, x2, out1, - out2): + for x1_i, x2_i, out1_i, out2_i in zip(self.elem, x2, + out1_seq, out2_seq): r1, r2 = getattr(x1_i.ufuncs, name)(x2_i, out1=out1_i, out2=out2_i, **kwargs) res1.append(r1) res2.append(r2) - out_space_1 = ProductSpace(*[r.space for r in res1]) - out_space_2 = ProductSpace(*[r.space for r in res2]) - return out_space_1.element(res1), out_space_2.element(res2) + + if out1 is None: + out_space_1 = ProductSpace(*[r.space for r in res1]) + out1 = out_space_1.element(res1) + if out2 is None: + out_space_2 = ProductSpace(*[r.space for r in res2]) + out2 = out_space_2.element(res2) + + return out1, out2 else: raise NotImplementedError From cc95a04da9fb495443415cdafb4e84a78341bdba Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Wed, 29 Nov 2017 11:29:43 +0100 Subject: [PATCH 25/38] ENH: implement Numpy-style broadcasting in pspace ufuncs --- odl/set/space.py | 8 +- odl/test/space/pspace_test.py | 17 ++++ odl/util/ufuncs.py | 157 +++++++++++++++++++++++++++++++--- 3 files changed, 166 insertions(+), 16 deletions(-) diff --git a/odl/set/space.py b/odl/set/space.py index 90effed68fd..a59c226b3ef 100644 --- a/odl/set/space.py +++ b/odl/set/space.py @@ -374,10 +374,10 @@ def __pow__(self, shape): >>> r2 ** 4 ProductSpace(rn(2), 4) - Multiple powers work as expected: + Multiple powers work as expected (first entry is outermost power): - >>> r2 ** (4, 2) - ProductSpace(ProductSpace(rn(2), 4), 2) + >>> r2 ** (3, 4) + ProductSpace(ProductSpace(rn(2), 4), 3) """ from odl.space import ProductSpace @@ -387,7 +387,7 @@ def __pow__(self, shape): shape = tuple(shape) pspace = self - for n in shape: + for n in reversed(shape): pspace = ProductSpace(pspace, n) return pspace diff --git a/odl/test/space/pspace_test.py b/odl/test/space/pspace_test.py index a28e310ebf0..efbbba0097f 100644 --- a/odl/test/space/pspace_test.py +++ b/odl/test/space/pspace_test.py @@ -944,6 +944,23 @@ def test_ufuncs(): assert w is z assert all_almost_equal(z, [[5], [7, 9]]) + # Broadcasting + pow_space = odl.rn(4) ** (2, 3) # corresponds to rn((2, 3, 4)) + x = pow_space.one() + x_arr = np.ones((2, 3, 4)) + y = x.ufuncs.add([1, 2, 3, 4]) # bcast along axes 0 and 1 + y_arr = np.add(x_arr, [1, 2, 3, 4]) + assert all_almost_equal(y, y_arr) + y = x.ufuncs.add(np.ones((3, 4))) # bcast along axis 0 + y_arr = np.add(x_arr, np.ones((3, 4))) + assert all_almost_equal(y, y_arr) + y = x.ufuncs.add(np.ones((2, 1, 4))) # bcast along axis 1 + y_arr = np.add(x_arr, np.ones((2, 1, 4))) + assert all_almost_equal(y, y_arr) + y = x.ufuncs.add((odl.rn(4) ** 3).one()) # bcast along axis 0 + y_arr = np.add(x_arr, np.ones((3, 4))) + assert all_almost_equal(y, y_arr) + def test_reductions(): H = odl.ProductSpace(odl.rn(1), odl.rn(2)) diff --git a/odl/util/ufuncs.py b/odl/util/ufuncs.py index 5e211f2a5a7..2545d17d252 100644 --- a/odl/util/ufuncs.py +++ b/odl/util/ufuncs.py @@ -224,20 +224,140 @@ def max(self, axis=None, dtype=None, out=None, keepdims=False): # --- Wrappers for `ProductSpaceElement` --- # +def _add_leading_dims(x2, pspace): + """Add leading dimensions to an input of a binary product space ufunc. + + Parameters + ---------- + x2 + Input to a binary ufunc. + pspace : `ProductSpace` + Spac on which the ufunc is evaluated. + + Returns + ------- + x2_extra_dims + Variant of ``x2`` with extra leading dimensions appropriate for + ``pspace``. + + Examples + -------- + Product space elements, arrays and nested sequences can be rewrapped + such that they have leading dimensions of size 1 for broadcasting: + + >>> pspace = odl.rn(2) ** (3, 4) + >>> reshaped = _add_leading_dims([0, 0], pspace) + >>> np.shape([0, 0]) + (2,) + >>> reshaped + [[[0, 0]]] + >>> np.shape(reshaped) + (1, 1, 2) + >>> reshaped = _add_leading_dims(np.zeros(2), pspace) + >>> reshaped + array([[[ 0., 0.]]]) + >>> reshaped.shape + (1, 1, 2) + >>> reshaped = _add_leading_dims(odl.rn(2).zero(), pspace) + >>> reshaped # returning array for simplicity + array([[[ 0., 0.]]]) + >>> reshaped.shape + (1, 1, 2) + >>> reshaped = _add_leading_dims((odl.rn(2) ** 4).zero(), pspace) + >>> reshaped # Broadcasting along outer dimension of size 3 + ProductSpace(ProductSpace(rn(2), 4), 1).element([ + [ + [ 0., 0.], + [ 0., 0.], + [ 0., 0.], + [ 0., 0.] + ] + ]) + """ + from odl.space.pspace import ProductSpace, ProductSpaceElement + from odl.space.base_tensors import Tensor + + # Workaround for `shape` not using the base space shape of a + # power space + # TODO: remove when fixed, see + # https://github.com/odlgroup/odl/pull/1152 + num_levels_pspace = 0 + tmp_pspace = pspace + while True: + if isinstance(tmp_pspace, ProductSpace): + tmp_pspace = tmp_pspace[0] + num_levels_pspace += 1 + else: + base_ndim_pspace = len(getattr(tmp_pspace, 'shape', ())) + break + + pspace_ndim = num_levels_pspace + base_ndim_pspace + + if isinstance(x2, ProductSpaceElement): + # Workaround for `shape` not using the base space shape of a + # power space element + # TODO: remove when fixed, see + # https://github.com/odlgroup/odl/pull/1152 + num_levels_x2 = 0 + tmp_x2 = x2 + while True: + if isinstance(tmp_x2, ProductSpaceElement): + tmp_x2 = tmp_x2[0] + num_levels_x2 += 1 + else: + base_ndim_x2 = len(getattr(tmp_x2, 'shape', ())) + break + + x2_ndim = num_levels_x2 + base_ndim_x2 + + # TODO: replace by `pspace.ndim` and `x2.ndim` when fixed, see + # https://github.com/odlgroup/odl/pull/1152 + for _ in range(pspace_ndim - x2_ndim): + x2 = (x2.space ** 1).element([x2]) + + return x2 + + if isinstance(x2, Tensor): + # Work with array directly + x2 = x2.data + + if hasattr(x2, 'shape'): + # Some type of array + x2_ndim = len(x2.shape) + + slc = (None,) * (pspace_ndim - x2_ndim) + (slice(None),) * x2_ndim + x2 = x2[slc] + + # Downstream code will raise in case of bad shape + return x2 + + else: + # Array-like + x2_ndim = np.ndim(x2) + + for _ in range(pspace_ndim - x2_ndim): + x2 = [x2] + + # Downstream code will raise in case of bad shape + return x2 + + def wrap_ufunc_productspace(name, n_in, n_out, doc): """Return ufunc wrapper for `ProductSpaceUfuncs`.""" if n_in == 1: if n_out == 1: def wrapper(self, out=None, **kwargs): - from odl.space.pspace import ProductSpace + from odl.space.pspace import ProductSpace, ProductSpaceElement if out is None: out_seq = [None] * len(self.elem.space) else: + assert isinstance(out, ProductSpaceElement) out_seq = out res = [] for xi, out_i in zip(self.elem, out_seq): - r = getattr(xi.ufuncs, name)(out=out_i, **kwargs) + ufunc = getattr(xi.ufuncs, name) + r = ufunc(out=out_i, **kwargs) res.append(r) if out is None: @@ -248,21 +368,22 @@ def wrapper(self, out=None, **kwargs): elif n_out == 2: def wrapper(self, out1=None, out2=None, **kwargs): - from odl.space.pspace import ProductSpace + from odl.space.pspace import ProductSpace, ProductSpaceElement if out1 is None: out1_seq = [None] * len(self.elem.space) else: + assert isinstance(out1, ProductSpaceElement) out1_seq = out1 if out2 is None: out2_seq = [None] * len(self.elem.space) else: + assert isinstance(out2, ProductSpaceElement) out2_seq = out2 res1, res2 = [], [] for xi, out1_i, out2_i in zip(self.elem, out1_seq, out2_seq): - r1, r2 = getattr(xi.ufuncs, name)(out1=out1_i, - out2=out2_i, - **kwargs) + ufunc = getattr(xi.ufuncs, name) + r1, r2 = ufunc(out1=out1_i, out2=out2_i, **kwargs) res1.append(r1) res2.append(r2) @@ -281,15 +402,22 @@ def wrapper(self, out1=None, out2=None, **kwargs): elif n_in == 2: if n_out == 1: def wrapper(self, x2, out=None, **kwargs): - from odl.space.pspace import ProductSpace + from odl.space.pspace import ProductSpace, ProductSpaceElement if out is None: out_seq = [None] * len(self.elem.space) else: + assert isinstance(out, ProductSpaceElement) out_seq = out + # Implement broadcasting + x2 = _add_leading_dims(x2, self.elem.space) + if len(x2) == 1 and len(self.elem.space) != 1: + x2 = [x2[0]] * len(self.elem.space) + res = [] for x1_i, x2_i, out_i in zip(self.elem, x2, out_seq): - r = getattr(x1_i.ufuncs, name)(x2_i, out=out_i, **kwargs) + ufunc = getattr(x1_i.ufuncs, name) + r = ufunc(x2_i, out=out_i, **kwargs) res.append(r) if out is None: @@ -300,22 +428,27 @@ def wrapper(self, x2, out=None, **kwargs): elif n_out == 2: def wrapper(self, x2, out1=None, out2=None, **kwargs): - from odl.space.pspace import ProductSpace + from odl.space.pspace import ProductSpace, ProductSpaceElement if out1 is None: out1_seq = [None] * len(self.elem.space) else: + assert isinstance(out1, ProductSpaceElement) out1_seq = out1 if out2 is None: out2_seq = [None] * len(self.elem.space) else: out2_seq = out2 + # Implement broadcasting + x2 = _add_leading_dims(x2, self.elem.space) + if len(x2) == 1 and len(self.elem.space) != 1: + x2 = [x2[0]] * len(self.elem.space) + res1, res2 = [], [] for x1_i, x2_i, out1_i, out2_i in zip(self.elem, x2, out1_seq, out2_seq): - r1, r2 = getattr(x1_i.ufuncs, name)(x2_i, out1=out1_i, - out2=out2_i, - **kwargs) + ufunc = getattr(x1_i.ufuncs, name) + r1, r2 = ufunc(x2_i, out1=out1_i, out2=out2_i, **kwargs) res1.append(r1) res2.append(r2) From 3084f1dcdccb612918e5359efad25b1b3280f15e Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Wed, 29 Nov 2017 16:21:39 +0100 Subject: [PATCH 26/38] WIP: simplify tests and add utils for array modules - expose trivial `none_context` - add `xfail_if` utility for conditional xfail - replace `almost_equal` by `pytest.approx` everywhere - extend more lp_discr tests to other impls - expose `array_module`, `array_cls` and `as_numpy` tools for handling arrays with different impls --- odl/space/base_tensors.py | 27 +- odl/space/npy_tensors.py | 25 +- odl/test/discr/diff_ops_test.py | 6 + odl/test/discr/discr_ops_test.py | 2 +- odl/test/discr/lp_discr_test.py | 384 ++++++++++-------- .../largescale/trafos/fourier_slow_test.py | 2 +- odl/test/operator/oputils_test.py | 25 ++ odl/test/set/domain_test.py | 3 + .../solvers/functional/functional_test.py | 32 +- odl/test/space/pspace_test.py | 43 +- odl/test/space/tensors_test.py | 157 ++++--- odl/test/util/numerics_test.py | 6 +- odl/util/testutils.py | 45 +- odl/util/utility.py | 60 +++ 14 files changed, 452 insertions(+), 365 deletions(-) diff --git a/odl/space/base_tensors.py b/odl/space/base_tensors.py index c909e6b0b16..f88bd6336ee 100644 --- a/odl/space/base_tensors.py +++ b/odl/space/base_tensors.py @@ -9,7 +9,6 @@ """Base classes for implementations of tensor spaces.""" from __future__ import print_function, division, absolute_import -from builtins import object from numbers import Integral import numpy as np @@ -18,7 +17,8 @@ from odl.util import ( is_numeric_dtype, is_real_dtype, is_floating_dtype, is_real_floating_dtype, is_complex_floating_dtype, safe_int_conv, - array_str, dtype_str, signature_string, indent, writable_array) + array_str, dtype_str, signature_string, indent, writable_array, + none_context) from odl.util.ufuncs import TensorSpaceUfuncs from odl.util.utility import TYPE_MAP_R2C, TYPE_MAP_C2R @@ -815,26 +815,11 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): # --- Evaluate ufunc --- # - # Trivial context used to create a single code path for the ufunc - # evaluation. For `None` output parameter(s), this is used instead of - # `writable_array`. - class CtxNone(object): - """Trivial context manager class. - - When used as :: - - with CtxNone() as obj: - # do stuff with `obj` - - the returned ``obj`` is ``None``. - """ - __enter__ = __exit__ = lambda *_: None - if method == '__call__': if ufunc.nout == 1: # Make context for output (trivial one returns `None`) if out is None: - out_ctx = CtxNone() + out_ctx = none_context() else: out_ctx = writable_array(out, **array_kwargs) @@ -851,11 +836,11 @@ class CtxNone(object): if out1 is not None: out1_ctx = writable_array(out1, **array_kwargs) else: - out1_ctx = CtxNone() + out1_ctx = none_context() if out2 is not None: out2_ctx = writable_array(out2, **array_kwargs) else: - out2_ctx = CtxNone() + out2_ctx = none_context() # Evaluate ufunc with out1_ctx as out1_arr, out2_ctx as out2_arr: @@ -872,7 +857,7 @@ class CtxNone(object): else: # method != '__call__' # Make context for output (trivial one returns `None`) if out is None: - out_ctx = CtxNone() + out_ctx = none_context() else: out_ctx = writable_array(out, **array_kwargs) diff --git a/odl/space/npy_tensors.py b/odl/space/npy_tensors.py index f3a87751bc5..8a80bcc4b75 100644 --- a/odl/space/npy_tensors.py +++ b/odl/space/npy_tensors.py @@ -23,7 +23,7 @@ CustomInner, CustomNorm, CustomDist) from odl.util import ( dtype_str, signature_string, is_real_dtype, is_numeric_dtype, - writable_array, is_floating_dtype) + writable_array, is_floating_dtype, none_context) __all__ = ('NumpyTensorSpace',) @@ -1659,26 +1659,11 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): # --- Evaluate ufunc --- # - # Trivial context used to create a single code path for the ufunc - # evaluation. For `None` output parameter(s), this is used instead of - # `writable_array`. - class CtxNone(object): - """Trivial context manager class. - - When used as :: - - with CtxNone() as obj: - # do stuff with `obj` - - the returned ``obj`` is ``None``. - """ - __enter__ = __exit__ = lambda *_: None - if method == '__call__': if ufunc.nout == 1: # Make context for output (trivial one returns `None`) if out is None: - out_ctx = CtxNone() + out_ctx = none_context() else: out_ctx = writable_array(out, **array_kwargs) @@ -1706,11 +1691,11 @@ class CtxNone(object): if out1 is not None: out1_ctx = writable_array(out1, **array_kwargs) else: - out1_ctx = CtxNone() + out1_ctx = none_context() if out2 is not None: out2_ctx = writable_array(out2, **array_kwargs) else: - out2_ctx = CtxNone() + out2_ctx = none_context() # Evaluate ufunc with out1_ctx as out1_arr, out2_ctx as out2_arr: @@ -1736,7 +1721,7 @@ class CtxNone(object): else: # method != '__call__' # Make context for output (trivial one returns `None`) if out is None: - out_ctx = CtxNone() + out_ctx = none_context() else: out_ctx = writable_array(out, **array_kwargs) diff --git a/odl/test/discr/diff_ops_test.py b/odl/test/discr/diff_ops_test.py index 5b37889d499..90d0a910ccc 100644 --- a/odl/test/discr/diff_ops_test.py +++ b/odl/test/discr/diff_ops_test.py @@ -430,6 +430,12 @@ def test_divergence(space, method, padding): assert rhs != 0 assert lhs == pytest.approx(rhs, rel=dtype_tol(space.dtype)) + # Higher dimensional arrays + for ndim in range(1, 6): + # DiscreteLpElement + lin_size = 3 + space = odl.uniform_discr([0.] * ndim, [1.] * ndim, [lin_size] * ndim) + # --- Laplacian --- # diff --git a/odl/test/discr/discr_ops_test.py b/odl/test/discr/discr_ops_test.py index 061cccbdfdb..bf8c8ca958d 100644 --- a/odl/test/discr/discr_ops_test.py +++ b/odl/test/discr/discr_ops_test.py @@ -9,8 +9,8 @@ """Unit tests for `discr_ops`.""" from __future__ import division -import pytest import numpy as np +import pytest import odl from odl.discr.discr_ops import _SUPPORTED_RESIZE_PAD_MODES diff --git a/odl/test/discr/lp_discr_test.py b/odl/test/discr/lp_discr_test.py index f829d28cf9b..3737b116f41 100644 --- a/odl/test/discr/lp_discr_test.py +++ b/odl/test/discr/lp_discr_test.py @@ -16,8 +16,9 @@ from odl.space.base_tensors import TensorSpace from odl.space.npy_tensors import NumpyTensor from odl.space.weighting import ConstWeighting +from odl.util import array_module, as_numpy from odl.util.testutils import ( - all_equal, all_almost_equal, noise_elements, simple_fixture) + all_equal, all_almost_equal, noise_elements, simple_fixture, xfail_if) USE_ARRAY_UFUNCS_INTERFACE = (parse_version(np.__version__) >= @@ -107,7 +108,8 @@ def test_empty(): # --- uniform_discr --- # -def test_factory_dtypes(odl_tspace_impl): +def test_uniform_discr_dtypes(odl_tspace_impl): + """Test basic properties of spaces created by uniform_discr.""" impl = odl_tspace_impl real_float_dtypes = [np.float32, np.float64] nonfloat_dtypes = [np.int8, np.int16, np.int32, np.int64, @@ -198,40 +200,39 @@ def test_uniform_discr_init_complex(odl_tspace_impl): # --- DiscreteLp methods --- # -def test_discretelp_element(): +def test_discretelp_element(tspace_impl): """Test creation and membership of DiscreteLp elements.""" # Creation from scratch # 1D - discr = odl.uniform_discr(0, 1, 3) + discr = odl.uniform_discr(0, 1, 3, impl=tspace_impl) weight = 1.0 if exponent == float('inf') else discr.cell_volume - tspace = odl.rn(3, weighting=weight) + tspace = odl.rn(3, weighting=weight, impl=tspace_impl) elem = discr.element() assert elem in discr assert elem.tensor in tspace # 2D - discr = odl.uniform_discr([0, 0], [1, 1], (3, 3)) + discr = odl.uniform_discr([0, 0], [1, 1], (3, 3), impl=tspace_impl) weight = 1.0 if exponent == float('inf') else discr.cell_volume - tspace = odl.rn((3, 3), weighting=weight) + tspace = odl.rn((3, 3), weighting=weight, impl=tspace_impl) elem = discr.element() assert elem in discr assert elem.tensor in tspace -def test_discretelp_element_from_array(): +def test_discretelp_element_from_array(tspace_impl): """Test creation of DiscreteLp elements from arrays.""" # 1D - discr = odl.uniform_discr(0, 1, 3) + discr = odl.uniform_discr(0, 1, 3, impl=tspace_impl) elem = discr.element([1, 2, 3]) assert np.array_equal(elem.tensor, [1, 2, 3]) - assert isinstance(elem, DiscreteLpElement) - assert isinstance(elem.tensor, NumpyTensor) assert all_equal(elem.tensor, [1, 2, 3]) -def test_element_from_array_2d(odl_elem_order): +def test_element_from_array_2d(odl_tspace_impl, odl_elem_order): """Test element in 2d with different orderings.""" + impl = odl_tspace_impl order = odl_elem_order discr = odl.uniform_discr([0, 0], [1, 1], [2, 2]) elem = discr.element([[1, 2], @@ -258,9 +259,9 @@ def test_element_from_array_2d(odl_elem_order): [4]]) # wrong shape -def test_element_from_function_1d(): +def test_element_from_function_1d(tspace_impl): """Test creation of DiscreteLp elements from functions in 1 dimension.""" - space = odl.uniform_discr(-1, 1, 4) + space = odl.uniform_discr(-1, 1, 4, impl=tspace_impl) points = space.points().squeeze() # Without parameter @@ -307,9 +308,9 @@ def f(x, **kwargs): assert all_equal(elem_lam, points) -def test_element_from_function_2d(): +def test_element_from_function_2d(tspace_impl): """Test creation of DiscreteLp elements from functions in 2 dimensions.""" - space = odl.uniform_discr([-1, -1], [1, 1], (2, 3)) + space = odl.uniform_discr([-1, -1], [1, 1], (2, 3), impl=tspace_impl) points = space.points() # Without parameter @@ -366,9 +367,9 @@ def f(x, **kwargs): assert all_equal(elem_lam, true_elem) -def test_discretelp_zero_one(): +def test_discretelp_zero_one(tspace_impl): """Test the zero and one element creators of DiscreteLp.""" - discr = odl.uniform_discr(0, 1, 3) + discr = odl.uniform_discr(0, 1, 3, impl=tspace_impl) zero = discr.zero() assert zero in discr @@ -742,12 +743,16 @@ def test_cell_volume(): assert elem.cell_volume == 0.5 -def test_astype(): +def test_astype(tspace_impl): - rdiscr = odl.uniform_discr([0, 0], [1, 1], [2, 2], dtype='float64') - cdiscr = odl.uniform_discr([0, 0], [1, 1], [2, 2], dtype='complex128') - rdiscr_s = odl.uniform_discr([0, 0], [1, 1], [2, 2], dtype='float32') - cdiscr_s = odl.uniform_discr([0, 0], [1, 1], [2, 2], dtype='complex64') + rdiscr = odl.uniform_discr([0, 0], [1, 1], [2, 2], dtype='float64', + impl=tspace_impl) + cdiscr = odl.uniform_discr([0, 0], [1, 1], [2, 2], dtype='complex128', + impl=tspace_impl) + rdiscr_s = odl.uniform_discr([0, 0], [1, 1], [2, 2], dtype='float32', + impl=tspace_impl) + cdiscr_s = odl.uniform_discr([0, 0], [1, 1], [2, 2], dtype='complex64', + impl=tspace_impl) # Real assert rdiscr.astype('float32') == rdiscr_s @@ -781,174 +786,202 @@ def testodl_ufuncs(odl_tspace_impl, odl_ufunc): space = odl.uniform_discr([0, 0], [1, 1], (2, 3), impl=impl) name = odl_ufunc - # Get the ufunc from numpy as reference - npy_ufunc = getattr(np, name) - nin = npy_ufunc.nin - nout = npy_ufunc.nout - if (np.issubsctype(space.dtype, np.floating) and - name in ['bitwise_and', - 'bitwise_or', - 'bitwise_xor', - 'invert', - 'left_shift', - 'right_shift']): - # Skip integer only methods if floating point type - return + # Get the ufunc from numpy as reference, plus some additional info + ufunc_npy = getattr(np, name) + nin = ufunc_npy.nin + nout = ufunc_npy.nout + + # Disable Numpy warnings for the time being + npy_err_orig = np.geterr() + np.seterr(all='ignore') + + def _check_result_type(result, expected_type): + if nout == 1: + assert isinstance(result, space.element_type) + elif nout > 1: + for i in range(nout): + assert isinstance(result[i], space.element_type) + else: + assert False + + def _check_result_is_out(result, out_seq): + if nout == 1: + assert result is out_seq[0] + elif nout > 1: + for i in range(nout): + assert result[i] is out_seq[i] + else: + assert False + + # See https://github.com/cupy/cupy/issues/794 + cupy_ufuncs_broken_complex = [ + 'expm1', 'floor_divide', 'fmin', 'fmax', + 'greater', 'greater_equal', 'less', 'less_equal', 'log1p', 'log2', + 'logical_and', 'logical_or', 'logical_not', 'logical_xor', 'minimum', + 'maximum', 'rint', 'sign', 'square'] + if (space.impl == 'cupy' and + space.dtype.kind == 'c' and + ufunc in cupy_ufuncs_broken_complex): + pytest.xfail('ufunc {} broken for complex input in cupy'.format(ufunc)) # Create some data arrays, elements = noise_elements(space, nin + nout) - in_arrays = arrays[:nin] - out_arrays = arrays[nin:] + # Arrays of the space's own data storage type + in_arrays_own = arrays[:nin] + in_arrays_npy = [as_numpy(arr) for arr in arrays[:nin]] data_elem = elements[0] - out_elems = elements[nin:] - - if nout == 1: - out_arr_kwargs = {'out': out_arrays[0]} - out_elem_kwargs = {'out': out_elems[0]} - elif nout > 1: - out_arr_kwargs = {'out': out_arrays[:nout]} - out_elem_kwargs = {'out': out_elems[:nout]} # Get function to call, using both interfaces: - # - vec.ufunc(other_args) - # - np.ufunc(vec, other_args) - elem_fun_old = getattr(data_elem.ufuncs, name) - in_elems_old = elements[1:nin] - elem_fun_new = npy_ufunc - in_elems_new = elements[:nin] - - # Out-of-place - with np.errstate(all='ignore'): # avoid pytest warnings - npy_result = npy_ufunc(*in_arrays) - odl_result_old = elem_fun_old(*in_elems_old) - assert all_almost_equal(npy_result, odl_result_old) - odl_result_new = elem_fun_new(*in_elems_new) - assert all_almost_equal(npy_result, odl_result_new) - - # Test type of output + # - vec.ufunc(*other_args) + # - np.ufunc(vec, *other_args) + ufunc_method = getattr(data_elem.ufuncs, name) + in_elems_method = elements[1:nin] + in_elems_npy = elements[:nin] + + # If Numpy fails, mark the test as xfail (same below) + try: + result_npy = ufunc_npy(*in_arrays_npy) + except TypeError: + pytest.xfail('numpy ufunc not valid for inputs') + + # Out-of-place -- in = elements -- ufunc = method or numpy + result = ufunc_method(*in_elems_method) + assert all_almost_equal(result_npy, result) + _check_result_type(result, space.element_type) + + # Get element(s) in the right space for in-place later if nout == 1: - assert isinstance(odl_result_old, space.element_type) - assert isinstance(odl_result_new, space.element_type) - elif nout > 1: - for i in range(nout): - assert isinstance(odl_result_old[i], space.element_type) - assert isinstance(odl_result_new[i], space.element_type) - - # In-place with ODL objects as `out` - with np.errstate(all='ignore'): # avoid pytest warnings - npy_result = npy_ufunc(*in_arrays, **out_arr_kwargs) - odl_result_old = elem_fun_old(*in_elems_old, **out_elem_kwargs) - assert all_almost_equal(npy_result, odl_result_old) - if USE_ARRAY_UFUNCS_INTERFACE: - # In-place will not work with Numpy < 1.13 - odl_result_new = elem_fun_new(*in_elems_new, **out_elem_kwargs) - assert all_almost_equal(npy_result, odl_result_new) - - # Check that returned stuff refers to given out + out_elems = [result.space.element()] + elif nout == 2: + out_elems = [res.space.element() for res in result] + else: + assert False + + result = ufunc_npy(*in_elems_npy) + assert all_almost_equal(result_npy, result) + _check_result_type(result, space.element_type) + + # Out-of-place -- in = numpy or own arrays -- ufunc = method + result = ufunc_method(*in_arrays_npy[1:]) + assert all_almost_equal(result_npy, result) + _check_result_type(result, space.element_type) + + result = ufunc_method(*in_arrays_own[1:]) + assert all_almost_equal(result_npy, result) + _check_result_type(result, space.element_type) + + # In-place -- in = elements -- out = elements -- ufunc = method or numpy if nout == 1: - assert odl_result_old is out_elems[0] - if USE_ARRAY_UFUNCS_INTERFACE: - assert odl_result_new is out_elems[0] - elif nout > 1: - for i in range(nout): - assert odl_result_old[i] is out_elems[i] - if USE_ARRAY_UFUNCS_INTERFACE: - assert odl_result_new[i] is out_elems[i] - - # In-place with Numpy array as `out` for new interface - if USE_ARRAY_UFUNCS_INTERFACE: - out_arrays_new = tuple(np.empty_like(arr) for arr in out_arrays) - if nout == 1: - out_arr_kwargs_new = {'out': out_arrays_new[0]} - elif nout > 1: - out_arr_kwargs_new = {'out': out_arrays_new[:nout]} + kwargs_out = {'out': out_elems[0]} + elif nout == 2: + kwargs_out = {'out1': out_elems[0], 'out2': out_elems[1]} - with np.errstate(all='ignore'): # avoid pytest warnings - odl_result_arr_new = elem_fun_new(*in_elems_new, - **out_arr_kwargs_new) - assert all_almost_equal(npy_result, odl_result_arr_new) + result = ufunc_method(*in_elems_method, **kwargs_out) + assert all_almost_equal(result_npy, result) + _check_result_is_out(result, out_elems[:nout]) + if USE_ARRAY_UFUNCS_INTERFACE: + # Custom objects not allowed as `out` for numpy < 1.13 if nout == 1: - assert odl_result_arr_new is out_arrays_new[0] - elif nout > 1: - for i in range(nout): - assert odl_result_arr_new[i] is out_arrays_new[i] + kwargs_out = {'out': out_elems[0]} + elif nout == 2: + kwargs_out = {'out': (out_elems[0], out_elems[1])} + + result = ufunc_npy(*in_elems_npy, **kwargs_out) + assert all_almost_equal(result_npy, result) + _check_result_is_out(result, out_elems[:nout]) - # In-place with data container (tensor) as `out` for new interface + # In-place -- in = elements -- out = numpy or own arrays -- ufunc = numpy + # This case is only supported with the new interface if USE_ARRAY_UFUNCS_INTERFACE: - out_tensors_new = tuple(space.tspace.element(np.empty_like(arr)) - for arr in out_arrays) + # Fresh arrays for output + out_arrays_npy = [np.empty_like(as_numpy(arr)) + for arr in arrays[nin:]] + out_arrays_own = [array_module(space.impl).empty_like(arr) + for arr in arrays[nin:]] if nout == 1: - out_tens_kwargs_new = {'out': out_tensors_new[0]} + kwargs_npy = {'out': out_arrays_npy[0]} + kwargs_own = {'out': out_arrays_own[0]} elif nout > 1: - out_tens_kwargs_new = {'out': out_tensors_new[:nout]} + kwargs_npy = {'out': out_arrays_npy[:nout]} + kwargs_own = {'out': out_arrays_own[:nout]} - with np.errstate(all='ignore'): # avoid pytest warnings - odl_result_tens_new = elem_fun_new(*in_elems_new, - **out_tens_kwargs_new) - assert all_almost_equal(npy_result, odl_result_tens_new) + try: + result_out_npy = ufunc_npy(*in_elems_npy, **kwargs_npy) + except TypeError: + pytest.xfail('numpy ufunc not valid for inputs') - if nout == 1: - assert odl_result_tens_new is out_tensors_new[0] - elif nout > 1: - for i in range(nout): - assert odl_result_tens_new[i] is out_tensors_new[i] + result_out_own = ufunc_npy(*in_elems_npy, **kwargs_own) + assert all_almost_equal(result_out_npy, result_npy) + assert all_almost_equal(result_out_own, result_npy) + _check_result_is_out(result_out_npy, out_arrays_npy) + _check_result_is_out(result_out_own, out_arrays_own) if USE_ARRAY_UFUNCS_INTERFACE: # Check `ufunc.at` indices = [[0, 0, 1], [0, 1, 2]] - mod_array = in_arrays[0].copy() - mod_elem = in_elems_new[0].copy() - if nout > 1: - return # currently not supported by Numpy + mod_array = in_arrays_npy[0].copy() + mod_elem = in_elems_npy[0].copy() if nin == 1: - with np.errstate(all='ignore'): # avoid pytest warnings - npy_result = npy_ufunc.at(mod_array, indices) - odl_result = npy_ufunc.at(mod_elem, indices) + try: + result_npy = ufunc_npy.at(mod_array, indices) + except TypeError: + pytest.xfail('numpy ufunc.at not valid for inputs') + + result = ufunc_npy.at(mod_elem, indices) + elif nin == 2: - other_array = in_arrays[1][indices] - other_elem = in_elems_new[1][indices] - with np.errstate(all='ignore'): # avoid pytest warnings - npy_result = npy_ufunc.at(mod_array, indices, other_array) - odl_result = npy_ufunc.at(mod_elem, indices, other_elem) + other_array = in_arrays_npy[1][indices] + other_elem = in_elems_npy[1][indices] + try: + result_npy = ufunc_npy.at(mod_array, indices, other_array) + except TypeError: + pytest.xfail('numpy ufunc.at not valid for inputs') + + result = ufunc_npy.at(mod_elem, indices, other_elem) - assert all_almost_equal(odl_result, npy_result) + assert all_almost_equal(result, result_npy) # Check `ufunc.reduce` if nin == 2 and nout == 1 and USE_ARRAY_UFUNCS_INTERFACE: - in_array = in_arrays[0] - in_elem = in_elems_new[0] + in_array = in_arrays_npy[0] + in_elem = in_elems_npy[0] # We only test along one axis since some binary ufuncs are not # re-orderable, in which case Numpy raises a ValueError - with np.errstate(all='ignore'): # avoid pytest warnings - npy_result = npy_ufunc.reduce(in_array) - odl_result = npy_ufunc.reduce(in_elem) - assert all_almost_equal(odl_result, npy_result) - # In-place using `out` (with ODL vector and array) - out_elem = odl_result.space.element() - out_array = np.empty(odl_result.shape, - dtype=odl_result.dtype) - npy_ufunc.reduce(in_elem, out=out_elem) - npy_ufunc.reduce(in_elem, out=out_array) - assert all_almost_equal(out_elem, odl_result) - assert all_almost_equal(out_array, odl_result) - # Using a specific dtype - try: - npy_result = npy_ufunc.reduce(in_array, dtype=complex) - except TypeError: - # Numpy finds no matching loop, bail out - return - else: - odl_result = npy_ufunc.reduce(in_elem, dtype=complex) - assert odl_result.dtype == npy_result.dtype - assert all_almost_equal(odl_result, npy_result) + + try: + result_npy = ufunc_npy.reduce(in_array) + except TypeError: + pytest.xfail('numpy ufunc.reduce not valid for inputs') + + # Out-of-place -- in = element + result = ufunc_npy.reduce(in_elem) + assert all_almost_equal(result, result_npy) + # keepdims not allowed + with pytest.raises(ValueError): + ufunc_npy.reduce(in_elem, keepdims=True) + + # Using a specific dtype + try: + result_npy = ufunc_npy.reduce(in_array, dtype=complex) + except TypeError: + pytest.xfail('numpy ufunc.reduce not valid for complex dtype') + + with xfail_if(space.impl == 'cupy', + reason='complex reduce is broken in cupy'): + result = ufunc_npy.reduce(in_elem, dtype=complex) + assert result.dtype == result_npy.dtype + assert all_almost_equal(result, result_npy) # Other ufunc method use the same interface, to we don't perform # extra tests for them. + # Reset Numpy err handling + np.seterr(**npy_err_orig) + def test_ufunc_corner_cases(odl_tspace_impl): """Check if some corner cases are handled correctly.""" @@ -972,10 +1005,12 @@ def test_ufunc_corner_cases(odl_tspace_impl): assert res.space == space # Check usage of `order` argument - for order in ('C', 'F'): - res = x.__array_ufunc__(np.sin, '__call__', x, order=order) - assert all_almost_equal(res, np.sin(x.asarray())) - assert res.tensor.data.flags[order + '_CONTIGUOUS'] + with xfail_if(tspace_impl == 'cupy', + reason='cupy does not accept `order` in ufuncs'): + for order in ('C', 'F'): + res = x.__array_ufunc__(np.sin, '__call__', x, order=order) + assert all_almost_equal(res, np.sin(x.asarray())) + assert res.tensor.data.flags[order + '_CONTIGUOUS'] # Check usage of `dtype` argument res = x.__array_ufunc__(np.sin, '__call__', x, dtype=complex) @@ -1125,16 +1160,20 @@ def test_real_imag(odl_tspace_impl, odl_elem_order): # With [:] assignment x = cdiscr.zero() - x.real[:] = assigntype([[2, 3], - [4, 5]]) - assert all_equal(x.real, [[2, 3], - [4, 5]]) + with xfail_if(tspace_impl == 'cupy', + reason='cupy x.real[:] does not mutate x'): + x.real[:] = assigntype([[2, 3], + [4, 5]]) + assert all_equal(x.real, [[2, 3], + [4, 5]]) x = cdiscr.zero() - x.imag[:] = assigntype([[2, 3], - [4, 5]]) - assert all_equal(x.imag, [[2, 3], - [4, 5]]) + with xfail_if(tspace_impl == 'cupy', + reason='cupy x.imag[:] does not mutate x'): + x.imag[:] = assigntype([[2, 3], + [4, 5]]) + assert all_equal(x.imag, [[2, 3], + [4, 5]]) # Setting with scalars x = cdiscr.zero() @@ -1163,7 +1202,8 @@ def test_reduction(odl_tspace_impl, odl_reduction): # Create some data x_arr, x = noise_elements(space, 1) - assert reduction(x_arr) == pytest.approx(getattr(x.ufuncs, name)()) + assert (reduction(as_numpy(x_arr)) == + pytest.approx(getattr(x.ufuncs, name)())) def test_power(odl_tspace_impl, power): @@ -1171,9 +1211,9 @@ def test_power(odl_tspace_impl, power): space = odl.uniform_discr([0, 0], [1, 1], [2, 2], impl=impl) x_arr, x = noise_elements(space, 1) - x_pos_arr = np.abs(x_arr) + x_pos_arr = array_module(tspace_impl).abs(x_arr) x_neg_arr = -x_pos_arr - x_pos = np.abs(x) + x_pos = x.ufuncs.absolute() x_neg = -x_pos if int(power) != power: @@ -1182,8 +1222,8 @@ def test_power(odl_tspace_impl, power): y += 0.1 with np.errstate(invalid='ignore'): - true_pos_pow = np.power(x_pos_arr, power) - true_neg_pow = np.power(x_neg_arr, power) + true_pos_pow = array_module(tspace_impl).power(x_pos_arr, power) + true_neg_pow = array_module(tspace_impl).power(x_neg_arr, power) if int(power) != power and impl == 'cuda': with pytest.raises(ValueError): diff --git a/odl/test/largescale/trafos/fourier_slow_test.py b/odl/test/largescale/trafos/fourier_slow_test.py index 9738b96e8a8..7ab5f78eecb 100644 --- a/odl/test/largescale/trafos/fourier_slow_test.py +++ b/odl/test/largescale/trafos/fourier_slow_test.py @@ -79,7 +79,7 @@ def charfun_freq_ball(x): ball_dom_ft = ft(ball_dom) ball_ran_ift = ft.adjoint(ball_ran) assert (ball_dom.inner(ball_ran_ift) == - pytest.approx(ball_ran.inner(ball_dom_ft), ndigits=1)) + pytest.approx(ball_ran.inner(ball_dom_ft), rel=0.1)) if __name__ == '__main__': diff --git a/odl/test/operator/oputils_test.py b/odl/test/operator/oputils_test.py index 54b0d63f0d5..e628681f407 100644 --- a/odl/test/operator/oputils_test.py +++ b/odl/test/operator/oputils_test.py @@ -78,6 +78,31 @@ def test_matrix_representation_product_to_product(): matrix representation will be ``(2, 3, 2, 3)``. """ n = 3 + rn = odl.rn(n) + A = np.random.rand(n, n) + Aop = odl.MatrixOperator(A) + + m = 2 + rm = odl.rn(m) + B = np.random.rand(m, m) + Bop = odl.MatrixOperator(B) + + ran_and_dom = ProductSpace(rn, rm) + + AB_matrix = np.vstack([np.hstack([A, np.zeros((n, m))]), + np.hstack([np.zeros((m, n)), B])]) + ABop = ProductSpaceOperator([[Aop, 0], + [0, Bop]], + ran_and_dom, ran_and_dom) + matrix_repr = matrix_representation(ABop) + + assert all_almost_equal(AB_matrix, matrix_repr) + + +def test_matrix_representation_product_to_product_two(): + # Verify that the matrix representation function returns the correct matrix + n = 3 + rn = odl.rn(n) A = np.random.rand(n, n) Aop = odl.MatrixOperator(A) diff --git a/odl/test/set/domain_test.py b/odl/test/set/domain_test.py index caa75590776..e0bd3453e6e 100644 --- a/odl/test/set/domain_test.py +++ b/odl/test/set/domain_test.py @@ -53,6 +53,9 @@ def test_min_pt(): set_ = IntervalProd([1], [2]) assert set_.min_pt == 1 + set_ = IntervalProd(1, 2) + assert set_.min_pt == 1 + set_ = IntervalProd([1, 2, 3], [5, 6, 7]) assert all_equal(set_.min_pt, [1, 2, 3]) diff --git a/odl/test/solvers/functional/functional_test.py b/odl/test/solvers/functional/functional_test.py index ca79f02bce6..0db26dc84af 100644 --- a/odl/test/solvers/functional/functional_test.py +++ b/odl/test/solvers/functional/functional_test.py @@ -355,12 +355,9 @@ def test_functional_sum(space): func1.gradient(x) + func2.gradient(x), ndigits) - assert ( - func_sum.derivative(x)(p) == - pytest.approx( - func1.gradient(x).inner(p) + func2.gradient(x).inner(p), - rel=rtol) - ) + assert (func_sum.derivative(x)(p) == + pytest.approx(func1.gradient(x).inner(p) + + func2.gradient(x).inner(p), rel=rtol)) # Verify that proximal raises with pytest.raises(NotImplementedError): @@ -396,10 +393,8 @@ def test_functional_plus_scalar(space): assert all_almost_equal(func_scalar_sum.gradient(x), func.gradient(x), ndigits) - assert ( - func_scalar_sum.derivative(x)(p) == - pytest.approx(func.gradient(x).inner(p), rel=rtol) - ) + assert (func_scalar_sum.derivative(x)(p) == + pytest.approx(func.gradient(x).inner(p), rel=rtol)) # Test proximal operator sigma = 1.2 @@ -407,11 +402,9 @@ def test_functional_plus_scalar(space): func.proximal(sigma)(x), ndigits) - # Test convex conjugate - assert ( - func_scalar_sum.convex_conj(x) == - pytest.approx(func.convex_conj(x) - scalar, rel=rtol) - ) + # Test convex conjugate functional + assert (func_scalar_sum.convex_conj(x) == + pytest.approx(func.convex_conj(x) - scalar, rel=rtol)) assert all_almost_equal(func_scalar_sum.convex_conj.gradient(x), func.convex_conj.gradient(x), @@ -472,11 +465,10 @@ def test_translation_of_functional(space): second_translation) # Evaluation - assert ( - double_translated_functional(x) == - pytest.approx(test_functional(x - translation - second_translation), - rel=rtol) - ) + assert (double_translated_functional(x) == + pytest.approx( + test_functional(x - translation - second_translation), + rel=rtol)) def test_translation_proximal_stepsizes(): diff --git a/odl/test/space/pspace_test.py b/odl/test/space/pspace_test.py index efbbba0097f..75b4b17358d 100644 --- a/odl/test/space/pspace_test.py +++ b/odl/test/space/pspace_test.py @@ -250,24 +250,29 @@ def test_metric(): HxH = odl.ProductSpace(H, H, exponent=1.0) w1 = HxH.element([v11, v12]) w2 = HxH.element([v21, v22]) - assert (HxH.dist(w1, w2) == - pytest.approx(H.dist(v11, v21) + H.dist(v12, v22))) + + dist = HxH.dist(w1, w2) + expected_dist = sum([H.dist(v11, v21), H.dist(v12, v22)]) + assert dist == pytest.approx(expected_dist) # 2-norm HxH = odl.ProductSpace(H, H, exponent=2.0) w1 = HxH.element([v11, v12]) w2 = HxH.element([v21, v22]) - assert ( - HxH.dist(w1, w2) == - pytest.approx((H.dist(v11, v21) ** 2 + H.dist(v12, v22) ** 2) ** 0.5) - ) + + dist = HxH.dist(w1, w2) + expected_dist = np.sqrt( + sum(d ** 2) for d in (H.dist(v11, v21), H.dist(v12, v22))) + assert dist == pytest.approx(expected_dist) # inf norm HxH = odl.ProductSpace(H, H, exponent=float('inf')) w1 = HxH.element([v11, v12]) w2 = HxH.element([v21, v22]) - assert (HxH.dist(w1, w2) == - pytest.approx(max(H.dist(v11, v21), H.dist(v12, v22)))) + + dist = HxH.dist(w1, w2) + expected_dist = max(H.dist(v11, v21), H.dist(v12, v22)) + assert dist == pytest.approx(expected_dist) def test_norm(): @@ -278,18 +283,27 @@ def test_norm(): # 1-norm HxH = odl.ProductSpace(H, H, exponent=1.0) w = HxH.element([v1, v2]) - assert HxH.norm(w) == pytest.approx(H.norm(v1) + H.norm(v2)) + + norm = HxH.norm(w) + expected_norm = sum([H.norm(v1), H.norm(v2)]) + assert norm == pytest.approx(expected_norm) # 2-norm HxH = odl.ProductSpace(H, H, exponent=2.0) w = HxH.element([v1, v2]) - assert (HxH.norm(w) == - pytest.approx((H.norm(v1) ** 2 + H.norm(v2) ** 2) ** (1 / 2.0))) + + norm = HxH.norm(w) + expected_norm = np.sqrt( + sum(n ** 2) for n in (H.norm(v1), H.norm(v2))) + assert norm == pytest.approx(expected_norm) # inf norm HxH = odl.ProductSpace(H, H, exponent=float('inf')) w = HxH.element([v1, v2]) - assert HxH.norm(w) == pytest.approx(max(H.norm(v1), H.norm(v2))) + + norm = HxH.norm(w) + expected_norm = max([H.norm(v1), H.norm(v2)]) + assert norm == pytest.approx(expected_norm) def test_inner(): @@ -303,7 +317,10 @@ def test_inner(): HxH = odl.ProductSpace(H, H) v = HxH.element([v1, v2]) u = HxH.element([u1, u2]) - assert HxH.inner(v, u) == pytest.approx(H.inner(v1, u1) + H.inner(v2, u2)) + + inner = HxH.inner(v, u) + expected_inner = sum([H.inner(v1, u1), H.inner(v2, u2)]) + assert inner == pytest.approx(expected_inner) def test_vector_weighting(exponent): diff --git a/odl/test/space/tensors_test.py b/odl/test/space/tensors_test.py index d1d1347b328..888ca2dc34c 100644 --- a/odl/test/space/tensors_test.py +++ b/odl/test/space/tensors_test.py @@ -29,9 +29,10 @@ CupyTensorSpaceCustomInner, CupyTensorSpaceCustomNorm, CupyTensorSpaceCustomDist, CUPY_AVAILABLE, cupy) +from odl.util import array_module, array_cls, as_numpy from odl.util.testutils import ( all_almost_equal, all_equal, simple_fixture, - noise_array, noise_element, noise_elements) + noise_array, noise_element, noise_elements, xfail_if) from odl.util.ufuncs import UFUNCS @@ -45,31 +46,6 @@ # Functions to return arrays, classes etc. corresponding to impls. Extend # when a new impl is available. -def _module(impl): - """Return the array module for ``impl``.""" - if impl == 'numpy': - return np - elif impl == 'cupy': - return cupy - else: - assert False - - -def _array_cls(impl): - """Return the array class for given impl.""" - return _module(impl).ndarray - - -def _as_numpy(array): - """Return a numpy.ndarray from the given array.""" - if isinstance(array, np.ndarray): - return array - elif isinstance(array, cupy.ndarray): - return cupy.asnumpy(array) - else: - assert False - - def _data_ptr(array): """Return the memory address of the given array (depending on impl).""" if isinstance(array, np.ndarray): @@ -82,7 +58,7 @@ def _data_ptr(array): def _pos_array(space): """Create an array with positive real entries for ``space``.""" - return _module(space.impl).abs(noise_array(space)) + 0.1 + return array_module(space.impl).abs(noise_array(space)) + 0.1 def _weighting_cls(impl, kind): @@ -272,7 +248,8 @@ def test_init_tspace_weighting(weight, exponent, odl_tspace_impl): # Need cast before using in space creation since # ArrayWeighting.__eq__ uses `arr1 is arr2` to check arrays weight = cupy.asarray(weight) - elif isinstance(weight, _array_cls(impl)): + + if isinstance(weight, array_cls(impl)): weighting_cls = _weighting_cls(impl, 'array') else: weighting_cls = _weighting_cls(impl, 'const') @@ -354,7 +331,7 @@ def test_element(tspace, odl_elem_order): assert elem.data.flags[order + '_CONTIGUOUS'] # From array (C order) - arr_c = _module(tspace.impl).ascontiguousarray(noise_array(tspace)) + arr_c = array_module(tspace.impl).ascontiguousarray(noise_array(tspace)) elem = tspace.element(arr_c, order=order) assert all_equal(elem, arr_c) assert elem.shape == elem.data.shape @@ -369,7 +346,7 @@ def test_element(tspace, odl_elem_order): assert elem.data.flags[order + '_CONTIGUOUS'] # From array (F order) - arr_f = _module(tspace.impl).asfortranarray(noise_array(tspace)) + arr_f = array_module(tspace.impl).asfortranarray(noise_array(tspace)) elem = tspace.element(arr_f, order=order) assert all_equal(elem, arr_f) assert elem.shape == elem.data.shape @@ -624,7 +601,8 @@ def test_power(tspace): """Test ``**`` against direct array exponentiation.""" [x_arr, y_arr], [x, y] = noise_elements(tspace, n=2) y_pos = tspace.element(y.ufuncs.absolute() + 0.1) - y_pos_arr = (_module(tspace.impl).abs(y_arr) + 0.1).astype(tspace.dtype) + y_pos_arr = array_module(tspace.impl).abs(y_arr) + 0.1 + y_pos_arr = y_pos_arr.astype(tspace.dtype) # Testing standard positive integer power out-of-place and in-place assert all_almost_equal(x ** 2, x_arr ** 2) @@ -769,7 +747,7 @@ def test_inner_exceptions(tspace): def test_norm(tspace): """Test the norm method against numpy.linalg.norm.""" xarr, x = noise_elements(tspace) - correct_norm = np.linalg.norm(_as_numpy(xarr.ravel())) + correct_norm = np.linalg.norm(as_numpy(xarr.ravel())) # Allow some error for single and half precision assert tspace.norm(x) == pytest.approx(correct_norm, rel=1e-2) @@ -795,8 +773,7 @@ def test_pnorm(exponent, odl_tspace_impl): for tspace in spaces: xarr, x = noise_elements(tspace) - correct_norm = _module(impl).linalg.norm( - xarr.ravel(), ord=exponent) + correct_norm = np.linalg.norm(as_numpy(xarr.ravel()), ord=exponent) assert tspace.norm(x) == pytest.approx(correct_norm) assert x.norm() == pytest.approx(correct_norm) @@ -805,7 +782,7 @@ def test_pnorm(exponent, odl_tspace_impl): def test_dist(tspace): """Test the dist method against numpy.linalg.norm of the difference.""" [xarr, yarr], [x, y] = noise_elements(tspace, n=2) - correct_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel())) + correct_dist = np.linalg.norm(as_numpy((xarr - yarr).ravel())) # Allow some error for single and half precision assert tspace.dist(x, y) == pytest.approx(correct_dist, rel=1e-2) @@ -836,9 +813,8 @@ def test_pdist(odl_tspace_impl, exponent): for space in spaces: [xarr, yarr], [x, y] = noise_elements(space, n=2) - correct_dist = _module(impl).linalg.norm( - (xarr - yarr).ravel(), ord=exponent) - + correct_dist = np.linalg.norm(as_numpy((xarr - yarr).ravel()), + ord=exponent) assert space.dist(x, y) == pytest.approx(correct_dist) assert x.dist(y) == pytest.approx(correct_dist) @@ -889,14 +865,14 @@ def test_element_setitem(odl_tspace_impl, setitem_indices): assert all_equal(x, x_arr) # Setting values with arrays - rhs_arr = _module(impl).ones(sliced_shape) + rhs_arr = array_module(tspace_impl).ones(sliced_shape) x_arr[setitem_indices] = rhs_arr x[setitem_indices] = rhs_arr assert all_equal(x, x_arr) # Setting values with a list of lists rhs_list = (-np.ones(sliced_shape)).tolist() - x_arr = _as_numpy(x_arr) + x_arr = as_numpy(x_arr) x_arr[setitem_indices] = rhs_list x[setitem_indices] = rhs_list assert all_equal(x, x_arr) @@ -1116,7 +1092,7 @@ def test_array_wrap_method(odl_tspace_impl): space = odl.tensor_space((3, 4), dtype='float32', exponent=1, weighting=2, impl=impl) x_arr, x = noise_elements(space) - y_arr = _module(impl).sin(x_arr) + y_arr = array_module(impl).sin(x_arr) y = np.sin(x) # Should yield again an ODL tensor assert all_equal(y, y_arr) @@ -1150,8 +1126,8 @@ def test_array_weighting_init(odl_tspace_impl, exponent): weighting_arr = weighting_cls(weight_arr, exponent=exponent) weighting_elem = weighting_cls(weight_elem, exponent=exponent) - assert isinstance(weighting_arr.array, _array_cls(impl)) - assert isinstance(weighting_elem.array, _array_cls(impl)) + assert isinstance(weighting_arr.array, array_cls(impl)) + assert isinstance(weighting_elem.array, array_cls(impl)) def test_array_weighting_array_is_valid(odl_tspace_impl): @@ -1240,7 +1216,7 @@ def test_array_weighting_inner(tspace): weight_arr = _pos_array(tspace) weighting_cls = _weighting_cls(tspace.impl, 'array') weighting = weighting_cls(weight_arr) - true_inner = np.vdot(_as_numpy(yarr), _as_numpy(xarr * weight_arr)) + true_inner = np.vdot(as_numpy(yarr), as_numpy(xarr * weight_arr)) # Allow some error for single and half precision assert weighting.inner(x, y) == pytest.approx(true_inner, rel=1e-2) @@ -1258,10 +1234,10 @@ def test_array_weighting_norm(tspace, exponent): weighting_cls = _weighting_cls(tspace.impl, 'array') weighting = weighting_cls(weight_arr, exponent=exponent) if exponent == float('inf'): - true_norm = np.linalg.norm(_as_numpy(xarr.ravel()), ord=float('inf')) + true_norm = np.linalg.norm(as_numpy(xarr.ravel()), ord=float('inf')) else: true_norm = np.linalg.norm( - _as_numpy((weight_arr ** (1 / exponent) * xarr).ravel()), + as_numpy((weight_arr ** (1 / exponent) * xarr).ravel()), ord=exponent) # Allow some error for single and half precision @@ -1276,11 +1252,11 @@ def test_array_weighting_dist(tspace, exponent): weighting_cls = _weighting_cls(tspace.impl, 'array') weighting = weighting_cls(weight_arr, exponent=exponent) if exponent == float('inf'): - true_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel()), + true_dist = np.linalg.norm(as_numpy((xarr - yarr).ravel()), ord=float('inf')) else: true_dist = np.linalg.norm( - _as_numpy((weight_arr ** (1 / exponent) * (xarr - yarr)).ravel()), + as_numpy((weight_arr ** (1 / exponent) * (xarr - yarr)).ravel()), ord=exponent) # Allow some error for single and half precision @@ -1350,7 +1326,7 @@ def test_const_weighting_inner(tspace): constant = 1.5 weighting_cls = _weighting_cls(tspace.impl, 'const') weighting = weighting_cls(constant) - true_inner = constant * np.vdot(_as_numpy(yarr), _as_numpy(xarr)) + true_inner = constant * np.vdot(as_numpy(yarr), as_numpy(xarr)) # Allow some error for single and half precision assert weighting.inner(x, y) == pytest.approx(true_inner, rel=1e-2) @@ -1372,7 +1348,7 @@ def test_const_weighting_norm(tspace, exponent): factor = 1.0 else: factor = constant ** (1 / exponent) - true_norm = factor * np.linalg.norm(_as_numpy(xarr.ravel()), ord=exponent) + true_norm = factor * np.linalg.norm(as_numpy(xarr.ravel()), ord=exponent) # Allow some error for single and half precision assert weighting.norm(x) == pytest.approx(true_norm, rel=1e-2) @@ -1389,7 +1365,7 @@ def test_const_weighting_dist(tspace, exponent): factor = 1.0 else: factor = constant ** (1 / exponent) - true_dist = factor * np.linalg.norm(_as_numpy((xarr - yarr).ravel()), + true_dist = factor * np.linalg.norm(as_numpy((xarr - yarr).ravel()), ord=exponent) # Allow some error for single and half precision @@ -1401,20 +1377,20 @@ def test_custom_inner(tspace): [xarr, yarr], [x, y] = noise_elements(tspace, 2) def inner(x, y): - return np.vdot(_as_numpy(y.data), _as_numpy(x.data)) + return np.vdot(as_numpy(y.data), as_numpy(x.data)) weighting_cls = _weighting_cls(tspace.impl, 'inner') w = weighting_cls(inner) w_same = weighting_cls(inner) - w_other = weighting_cls(_module(tspace.impl).dot) + w_other = weighting_cls(array_module(tspace.impl).dot) assert w == w assert w == w_same assert w != w_other - true_inner = np.vdot(_as_numpy(yarr), _as_numpy(xarr)) - true_norm = np.linalg.norm(_as_numpy(xarr.ravel())) - true_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel())) + true_inner = np.vdot(as_numpy(yarr), as_numpy(xarr)) + true_norm = np.linalg.norm(as_numpy(xarr.ravel())) + true_dist = np.linalg.norm(as_numpy((xarr - yarr).ravel())) # Allow some error for single and half precision assert w.inner(x, y) == pytest.approx(true_inner, rel=1e-2) @@ -1430,10 +1406,10 @@ def test_custom_norm(tspace): [xarr, yarr], [x, y] = noise_elements(tspace, 2) def norm(x): - return np.linalg.norm(_as_numpy(x.data).ravel()) + return np.linalg.norm(as_numpy(x.data).ravel()) def other_norm(x): - return np.linalg.norm(_as_numpy(x.data).ravel(), ord=1) + return np.linalg.norm(as_numpy(x.data).ravel(), ord=1) weighting_cls = _weighting_cls(tspace.impl, 'norm') w = weighting_cls(norm) @@ -1444,8 +1420,8 @@ def other_norm(x): assert w == w_same assert w != w_other - true_norm = np.linalg.norm(_as_numpy(xarr.ravel())) - true_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel())) + true_norm = np.linalg.norm(as_numpy(xarr.ravel())) + true_dist = np.linalg.norm(as_numpy((xarr - yarr).ravel())) # Allow some error for single and half precision assert w.norm(x) == pytest.approx(true_norm, rel=1e-2) @@ -1463,10 +1439,10 @@ def test_custom_dist(tspace): [xarr, yarr], [x, y] = noise_elements(tspace, 2) def dist(x, y): - return np.linalg.norm(_as_numpy((x - y).data).ravel()) + return np.linalg.norm(as_numpy((x - y).data).ravel()) def other_dist(x, y): - return np.linalg.norm(_as_numpy((x - y).data).ravel(), ord=1) + return np.linalg.norm(as_numpy((x - y).data).ravel(), ord=1) weighting_cls = _weighting_cls(tspace.impl, 'dist') w = weighting_cls(dist) @@ -1477,7 +1453,7 @@ def other_dist(x, y): assert w == w_same assert w != w_other - true_dist = np.linalg.norm(_as_numpy((xarr - yarr).ravel())) + true_dist = np.linalg.norm(as_numpy((xarr - yarr).ravel())) # Allow some error for single and half precision assert w.dist(x, y) == pytest.approx(true_dist, rel=1e-2) @@ -1541,7 +1517,7 @@ def _check_result_is_out(result, out_seq): arrays, elements = noise_elements(tspace, nin + nout) # Arrays of the space's own data storage type in_arrays_own = arrays[:nin] - in_arrays_npy = [_as_numpy(arr) for arr in arrays[:nin]] + in_arrays_npy = [as_numpy(arr) for arr in arrays[:nin]] data_elem = elements[0] out_elems = elements[nin:] if nout == 1: @@ -1557,6 +1533,8 @@ def _check_result_is_out(result, out_seq): ufunc_method = getattr(data_elem.ufuncs, name) in_elems_method = elements[1:nin] in_elems_npy = elements[:nin] + + # If Numpy fails, mark the test as xfail (same below) try: result_npy = ufunc_npy(*in_arrays_npy) except TypeError: @@ -1611,9 +1589,9 @@ def _check_result_is_out(result, out_seq): # This case is only supported with the new interface if USE_ARRAY_UFUNCS_INTERFACE: # Fresh arrays for output - out_arrays_npy = [np.empty_like(_as_numpy(arr)) + out_arrays_npy = [np.empty_like(as_numpy(arr)) for arr in arrays[nin:]] - out_arrays_own = [_module(tspace.impl).empty_like(arr) + out_arrays_own = [array_module(tspace.impl).empty_like(arr) for arr in arrays[nin:]] if nout == 1: kwargs_npy = {'out': out_arrays_npy[0]} @@ -1687,7 +1665,7 @@ def _check_result_is_out(result, out_seq): dtype=result_keepdims.dtype) ufunc_npy.reduce(in_elem, out=out_array_npy, keepdims=True) assert all_almost_equal(out_array_npy, result_keepdims) - out_array_own = _module(tspace.impl).empty( + out_array_own = array_module(tspace.impl).empty( result_keepdims.shape, dtype=result_keepdims.dtype) ufunc_npy.reduce(in_elem, out=out_array_own, keepdims=True) assert all_almost_equal(out_array_own, result_keepdims) @@ -1698,12 +1676,10 @@ def _check_result_is_out(result, out_seq): except TypeError: pytest.xfail('numpy ufunc.reduce not valid for complex dtype') - if tspace.impl == 'cupy': - pytest.xfail('cupy ufunc.reduce raises error for complex dtype') - - result = ufunc_npy.reduce(in_elem, dtype=complex) - assert result.dtype == result_npy.dtype - assert all_almost_equal(result, result_npy) + with xfail_if(tspace.impl == 'cupy'): + result = ufunc_npy.reduce(in_elem, dtype=complex) + assert result.dtype == result_npy.dtype + assert all_almost_equal(result, result_npy) # Other ufunc method use the same interface, to we don't perform # extra tests for them. @@ -1770,6 +1746,8 @@ def test_ufunc_corner_cases(odl_tspace_impl): elif impl == 'cupy': with pytest.xfail(reason='cupy does not accept `order` in ufuncs'): res = x.__array_ufunc__(np.sin, '__call__', x, order=order) + assert all_almost_equal(res, np.sin(x.asarray())) + assert res.data.flags[order + '_CONTIGUOUS'] # Check usage of `dtype` argument res = x.__array_ufunc__(np.sin, '__call__', x, dtype='float32') @@ -1810,7 +1788,8 @@ def test_ufunc_corner_cases(odl_tspace_impl): res = x.__array_ufunc__(np.add, 'accumulate', x) assert all_almost_equal(res, np.add.accumulate(x.asarray())) assert res.space == space - arr = _module(impl).empty_like(x) + + arr = array_module(impl).empty_like(x) res = x.__array_ufunc__(np.add, 'accumulate', x, out=(arr,)) assert all_almost_equal(arr, np.add.accumulate(x.asarray())) assert res is arr @@ -1831,7 +1810,7 @@ def test_ufunc_corner_cases(odl_tspace_impl): assert all_almost_equal(res, np.add.reduce(x.asarray())) # With `out` argument and `axis` - out_ax0 = _module(impl).empty(3) + out_ax0 = array_module(impl).empty(3) res = x.__array_ufunc__(np.add, 'reduce', x, axis=0, out=(out_ax0,)) assert all_almost_equal(out_ax0, np.add.reduce(x.asarray(), axis=0)) assert res is out_ax0 @@ -1877,14 +1856,14 @@ def testodl_reduction(tspace, odl_reduction): pytest.xfail('Cupy does not accept complex input to `min` and `max`') # Full reduction, produces scalar - result_npy = npy_reduction(_as_numpy(x_arr)) + result_npy = npy_reduction(as_numpy(x_arr)) result = x_reduction() assert result == pytest.approx(result_npy) result = x_reduction(axis=(0, 1)) assert result == pytest.approx(result_npy) # Reduction along axes, produces element in reduced space - result_npy = npy_reduction(_as_numpy(x_arr), axis=0) + result_npy = npy_reduction(as_numpy(x_arr), axis=0) result = x_reduction(axis=0) assert isinstance(result, tspace.element_type) assert result.shape == result_npy.shape @@ -1900,7 +1879,7 @@ def testodl_reduction(tspace, odl_reduction): assert all_almost_equal(out, result_npy) # Use keepdims parameter - result_npy = npy_reduction(_as_numpy(x_arr), axis=1, keepdims=True) + result_npy = npy_reduction(as_numpy(x_arr), axis=1, keepdims=True) result = x_reduction(axis=1, keepdims=True) assert result.shape == result_npy.shape assert all_almost_equal(result, result_npy) @@ -1913,18 +1892,18 @@ def testodl_reduction(tspace, odl_reduction): # These reductions have a `dtype` parameter if name in ('cumprod', 'cumsum', 'mean', 'prod', 'std', 'sum', 'trace', 'var'): - if tspace.impl == 'cupy' and tspace.dtype == 'float16': + with xfail_if(tspace.impl == 'cupy' and tspace.dtype == 'float16', + reason='complex reduction fails for float16 in cupy'): # See https://github.com/cupy/cupy/issues/795 - pytest.xfail('reduction with complex fails for float16 in cupy') - - result_npy = npy_reduction(_as_numpy(x_arr), axis=1, dtype='complex64') - result = x_reduction(axis=1, dtype='complex64') - assert result.dtype == np.dtype('complex64') - assert all_almost_equal(result, result_npy) - # Evaluate in-place - out = result.space.element() - x_reduction(axis=1, dtype='complex64', out=out) - assert all_almost_equal(out, result_npy) + result_npy = npy_reduction(as_numpy(x_arr), axis=1, + dtype='complex64') + result = x_reduction(axis=1, dtype='complex64') + assert result.dtype == np.dtype('complex64') + assert all_almost_equal(result, result_npy) + # Evaluate in-place + out = result.space.element() + x_reduction(axis=1, dtype='complex64', out=out) + assert all_almost_equal(out, result_npy) def test_ufunc_reduction_docs_notempty(odl_tspace_impl): diff --git a/odl/test/util/numerics_test.py b/odl/test/util/numerics_test.py index a4a3fec702a..1fb8269464b 100644 --- a/odl/test/util/numerics_test.py +++ b/odl/test/util/numerics_test.py @@ -355,9 +355,9 @@ def test_resize_array_adj(resize_setup, odl_floating_dtype): resized_adj = resize_array(other_arr, array.shape, offset, pad_mode, pad_const, direction='adjoint') - assert (np.vdot(resized.ravel(), other_arr.ravel()) == - pytest.approx(np.vdot(array.ravel(), resized_adj.ravel()), - rel=dtype_tol(dtype))) + dot = np.vdot(resized.ravel(), other_arr.ravel()) + expected_dot = np.vdot(array.ravel(), resized_adj.ravel()) + assert dot == pytest.approx(expected_dot) def test_resize_array_corner_cases(odl_scalar_dtype, padding): diff --git a/odl/util/testutils.py b/odl/util/testutils.py index 808aaa5b2e5..10d9d97f44f 100644 --- a/odl/util/testutils.py +++ b/odl/util/testutils.py @@ -17,7 +17,7 @@ import warnings from time import time -from odl.util.utility import run_from_ipython, is_string +from odl.util.utility import run_from_ipython, is_string, none_context __all__ = ( @@ -208,6 +208,14 @@ def is_subdict(subdict, dictionary): return all(item in dictionary.items() for item in subdict.items()) +def xfail_if(condition, reason=''): + """Return a ``pytest.xfail`` object if ``condition`` is ``True``.""" + if condition: + return pytest.xfail(reason) + else: + return none_context() + + try: # Try catch in case user does not have pytest import pytest @@ -353,20 +361,12 @@ def noise_array(space): from odl.space.cupy_tensors import cupy if isinstance(space, ProductSpace): - arr_list = [noise_array(si) for si in space] - try: - impl = space[0].impl - except (IndexError, AttributeError): - impl = 'numpy' - + arr_list = [noise_array(spc_i) for spc_i in space] if space.is_power_space: - if impl == 'numpy': - return np.vstack(arr_list) - elif impl == 'cupy': - return cupy.vstack(arr_list) - else: - raise RuntimeError('bad `impl` {!r}'.format(impl)) - + arr = np.empty((len(arr_list),) + arr_list[0].shape, + dtype=space.dtype) + for i in range(len(arr)): + arr[i] = arr_list[i] else: return tuple(arr_list) @@ -386,14 +386,7 @@ def noise_array(space): else: raise ValueError('bad dtype {}'.format(space.dtype)) - arr = arr.astype(space.dtype, copy=False) - impl = getattr(space, 'impl', 'numpy') - if impl == 'numpy': - return arr - elif impl == 'cupy': - return cupy.asarray(arr) - else: - raise RuntimeError('bad `impl` {!r}'.format(impl)) + return arr.astype(space.dtype, copy=False) def noise_element(space): @@ -486,10 +479,12 @@ def noise_elements(space, n=1): noise_array noise_element """ - arrs = tuple(noise_array(space) for _ in range(n)) + npy_arrs = [noise_array(space) for _ in range(n)] - # Make space elements from arrays - elems = tuple(space.element(arr.copy()) for arr in arrs) + # Make space elements from Numpy arrays + elems = tuple(space.element(arr) for arr in npy_arrs) + # Make copies of the arrays + arrs = tuple(elem.data.copy() for elem in elems) if n == 1: return tuple(arrs + elems) diff --git a/odl/util/utility.py b/odl/util/utility.py index ce23fe14dc8..80928ba3252 100644 --- a/odl/util/utility.py +++ b/odl/util/utility.py @@ -31,6 +31,14 @@ 'cache_arguments', 'unique', 'REPR_PRECISION') +__all__ = ('array_str', 'dtype_str', 'dtype_repr', 'npy_printoptions', + 'signature_string', 'indent', + 'is_numeric_dtype', 'is_int_dtype', 'is_floating_dtype', + 'is_real_dtype', 'is_real_floating_dtype', + 'is_complex_floating_dtype', 'real_dtype', 'complex_dtype', + 'is_string', 'conj_exponent', 'writable_array', 'none_context', + 'array_module', 'array_cls', 'as_numpy', + 'run_from_ipython', 'NumpyRandomSeed', 'cache_arguments', 'unique') REPR_PRECISION = 4 # For printing scalars and array entries TYPE_MAP_R2C = {np.dtype(dtype): np.result_type(dtype, 1j) @@ -729,6 +737,58 @@ def __exit__(self, type, value, traceback): self.arr = None +class none_context(object): + """Trivial context manager class. + + When used as :: + + with none_context(*args, **kwargs) as obj: + # do stuff with `obj` + + the returned ``obj`` is ``None``. + """ + def __init__(self, *args, **kwargs): + # Ignore all arguments + pass + + def __enter__(self): + return None + + def __exit__(self, type, value, traceback): + return None + + +def array_module(impl): + """Return the array module for ``impl``.""" + from odl.space.cupy_tensors import cupy + if impl == 'numpy': + return np + elif impl == 'cupy': + return cupy + else: + raise ValueError('`impl` {!r} not understood'.format(impl)) + + +def array_cls(impl): + """Return the array class for given ``impl``.""" + return array_module(impl).ndarray + + +def as_numpy(array): + """Return a ``numpy.ndarray`` from the given array. + + This is intended for casting from other array implementations to + Numpy, not for ODL tensors. + """ + from odl.space.cupy_tensors import cupy + if isinstance(array, np.ndarray): + return array + elif isinstance(array, cupy.ndarray): + return cupy.asnumpy(array) + else: + raise TypeError('type {} not understood'.format(type(array))) + + def signature_string(posargs, optargs, sep=', ', mod='!r'): """Return a stringified signature from given arguments. From e3e364b7a968160309e60aa4c2959d48fa2b6d7b Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Wed, 29 Nov 2017 17:50:55 +0100 Subject: [PATCH 27/38] ENH: add impl to writable_array --- odl/util/utility.py | 44 +++++++++++++++++++++++++++++++------------- 1 file changed, 31 insertions(+), 13 deletions(-) diff --git a/odl/util/utility.py b/odl/util/utility.py index 80928ba3252..b91733c4fdb 100644 --- a/odl/util/utility.py +++ b/odl/util/utility.py @@ -657,17 +657,21 @@ def ip_wrapper(x, out, **kwargs): class writable_array(object): + """Context manager that casts obj to a `numpy.array` and saves changes.""" - def __init__(self, obj, **kwargs): - """initialize a new instance. + def __init__(self, obj, impl='numpy', **kwargs): + """Initialize a new instance. Parameters ---------- obj : `array-like` Object that should be made available as writable array. It must be valid as input to `numpy.asarray` and needs to - support the syntax ``obj[:] = arr``. + support assignment ``obj[:] = arr``. + impl : str, optional + Array backend for the exposed ndarray. + kwargs : Keyword arguments that should be passed to `numpy.asarray`. @@ -710,31 +714,45 @@ def __init__(self, obj, **kwargs): """ self.obj = obj self.kwargs = kwargs + self.impl = impl self.arr = None def __enter__(self): - """called by ``with writable_array(obj):``. + """Called by ``with writable_array(obj):``. Returns ------- - arr : `numpy.ndarray` - Array representing ``self.obj``, created by calling - ``numpy.asarray``. Any changes to ``arr`` will be passed through - to ``self.obj`` after the context manager exits. + arr : ndarray + Array representing ``self.obj``, created by calling ``asarray`` + corresponding to the chosen ``impl``, e.g., ``numpy.asarray``. + Any changes to ``arr`` will be passed through to ``self.obj`` + when the context manager exits. """ - self.arr = np.asarray(self.obj, **self.kwargs) + self.arr = array_module(self.impl).asarray(self.obj, **self.kwargs) return self.arr - def __exit__(self, type, value, traceback): - """called when ``with writable_array(obj):`` ends. + def __exit__(self, *args, **kwargs): + """Called when ``with writable_array(obj):`` ends. Saves any changes to ``self.arr`` to ``self.obj``, also "frees" self.arr in case the manager is used multiple times. Extra arguments are ignored, any exceptions are passed through. """ - self.obj[:] = self.arr - self.arr = None + # Some extra care for Numpy and Cupy arrays + from odl.space.npy_tensors import NumpyTensor + from odl.space.cupy_tensors import cupy + + obj_is_npy = isinstance(self.obj, (NumpyTensor, np.ndarray)) + arr_is_cupy = isinstance(self.arr, + getattr(cupy, 'ndarray', type(None))) + if obj_is_npy and arr_is_cupy: + arr = cupy.asnumpy(self.arr) + else: + arr = self.arr + + self.obj[:] = arr + self.arr = None # gc class none_context(object): From d0c970554c8675fc81de1695e968b261611bc7d6 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Wed, 29 Nov 2017 18:57:25 +0100 Subject: [PATCH 28/38] ENH: add impl to asarray, tests, and fix real and imag for cupy - use native impl in diff_ops - remove CupyTensor.ufuncs, use base class method - add skip_if_no_cupy decorator --- odl/discr/diff_ops.py | 65 ++++++++++-------- odl/discr/discretization.py | 14 ++-- odl/space/cupy_tensors.py | 117 +++++++++++++++------------------ odl/space/npy_tensors.py | 61 +++++++++++++---- odl/test/space/tensors_test.py | 82 +++++++++++++++++++++-- odl/util/testutils.py | 8 ++- 6 files changed, 230 insertions(+), 117 deletions(-) diff --git a/odl/discr/diff_ops.py b/odl/discr/diff_ops.py index fad6dbb939f..ba83586bf6b 100644 --- a/odl/discr/diff_ops.py +++ b/odl/discr/diff_ops.py @@ -14,7 +14,7 @@ from odl.discr.lp_discr import DiscreteLp from odl.operator.tensor_ops import PointwiseTensorFieldOperator from odl.space import ProductSpace -from odl.util import writable_array, signature_string, indent +from odl.util import writable_array, signature_string, indent, array_module __all__ = ('PartialDerivative', 'Gradient', 'Divergence', 'Laplacian') @@ -137,9 +137,9 @@ def _call(self, x, out=None): if out is None: out = self.range.element() - # TODO: this pipes CUDA arrays through NumPy. Write native operator. - with writable_array(out) as out_arr: - finite_diff(x.asarray(), axis=self.axis, dx=self.dx, + impl = self.domain.impl + with writable_array(out, impl=impl) as out_arr: + finite_diff(x.asarray(), axis=self.axis, dx=self.dx, impl=impl, method=self.method, pad_mode=self.pad_mode, pad_const=self.pad_const, out=out_arr) return out @@ -350,13 +350,13 @@ def _call(self, x, out=None): x_arr = x.asarray() ndim = self.domain.ndim dx = self.domain.cell_sides + impl = self.domain.impl for axis in range(ndim): - with writable_array(out[axis]) as out_arr: - finite_diff(x_arr, axis=axis, dx=dx[axis], method=self.method, - pad_mode=self.pad_mode, - pad_const=self.pad_const, - out=out_arr) + with writable_array(out[axis], impl=impl) as out_arr: + finite_diff(x_arr, axis=axis, dx=dx[axis], impl=impl, + method=self.method, pad_mode=self.pad_mode, + pad_const=self.pad_const, out=out_arr) return out def derivative(self, point=None): @@ -559,11 +559,13 @@ def _call(self, x, out=None): ndim = self.range.ndim dx = self.range.cell_sides + impl = self.range.impl - tmp = np.empty(out.shape, out.dtype, order=out.space.default_order) - with writable_array(out) as out_arr: + tmp = array_module(impl).empty( + out.shape, out.dtype, order=out.space.default_order) + with writable_array(out, impl=impl) as out_arr: for axis in range(ndim): - finite_diff(x[axis], axis=axis, dx=dx[axis], + finite_diff(x[axis], axis=axis, dx=dx[axis], impl=impl, method=self.method, pad_mode=self.pad_mode, pad_const=self.pad_const, out=tmp) @@ -714,24 +716,26 @@ def _call(self, x, out=None): else: out.set_zero() - x_arr = x.asarray() - out_arr = out.asarray() - tmp = np.empty(out.shape, out.dtype, order=out.space.default_order) - ndim = self.domain.ndim dx = self.domain.cell_sides + impl = self.domain.impl + + x_arr = x.asarray(impl=impl) + out_arr = out.asarray(impl=impl) + tmp = array_module(impl).empty( + out.shape, out.dtype, order=out.space.default_order) - with writable_array(out) as out_arr: + with writable_array(out, impl=impl) as out_arr: for axis in range(ndim): # TODO: this can be optimized - finite_diff(x_arr, axis=axis, dx=dx[axis] ** 2, + finite_diff(x_arr, axis=axis, dx=dx[axis] ** 2, impl=impl, method='forward', pad_mode=self.pad_mode, pad_const=self.pad_const, out=tmp) out_arr += tmp - finite_diff(x_arr, axis=axis, dx=dx[axis] ** 2, + finite_diff(x_arr, axis=axis, dx=dx[axis] ** 2, impl=impl, method='backward', pad_mode=self.pad_mode, pad_const=self.pad_const, out=tmp) @@ -785,7 +789,8 @@ def __str__(self): return '{}:\n{}'.format(self.__class__.__name__, indent(dom_ran_str)) -def finite_diff(f, axis, dx=1.0, method='forward', out=None, **kwargs): +def finite_diff(f, axis, dx=1.0, method='forward', out=None, impl='numpy', + **kwargs): """Calculate the partial derivative of ``f`` along a given ``axis``. In the interior of the domain of f, the partial derivative is computed @@ -819,6 +824,9 @@ def finite_diff(f, axis, dx=1.0, method='forward', out=None, **kwargs): out : `numpy.ndarray`, optional An N-dimensional array to which the output is written. Has to have the same shape as the input array ``f``. + impl : str, optional + Implementation backend for array manipulations. Usually handled by + the calling code. pad_mode : string, optional The padding mode to use outside the domain. @@ -883,7 +891,8 @@ def finite_diff(f, axis, dx=1.0, method='forward', out=None, **kwargs): >>> out is finite_diff(f, axis=0, out=out) True """ - f_arr = np.asarray(f) + arrmod = array_module(impl) + f_arr = arrmod.asarray(f) ndim = f_arr.ndim if f_arr.shape[axis] < 2: @@ -913,7 +922,7 @@ def finite_diff(f, axis, dx=1.0, method='forward', out=None, **kwargs): pad_const = f.dtype.type(pad_const) if out is None: - out = np.empty_like(f_arr) + out = arrmod.empty_like(f_arr) else: if out.shape != f.shape: raise ValueError('expected output shape {}, got {}' @@ -931,24 +940,24 @@ def finite_diff(f, axis, dx=1.0, method='forward', out=None, **kwargs): # create slice objects: initially all are [:, :, ..., :] - # Swap axes so that the axis of interest is first. This is a O(1) + # Swap axes so that the axis of interest is first. This is an O(1) # operation and is done to simplify the code below. - out, out_in = np.swapaxes(out, 0, axis), out - f_arr = np.swapaxes(f_arr, 0, axis) + out, out_in = arrmod.swapaxes(out, 0, axis), out + f_arr = arrmod.swapaxes(f_arr, 0, axis) # Interior of the domain of f if method == 'central': # 1D equivalent: out[1:-1] = (f[2:] - f[:-2])/2.0 - np.subtract(f_arr[2:], f_arr[:-2], out=out[1:-1]) + arrmod.subtract(f_arr[2:], f_arr[:-2], out=out[1:-1]) out[1:-1] /= 2.0 elif method == 'forward': # 1D equivalent: out[1:-1] = (f[2:] - f[1:-1]) - np.subtract(f_arr[2:], f_arr[1:-1], out=out[1:-1]) + arrmod.subtract(f_arr[2:], f_arr[1:-1], out=out[1:-1]) elif method == 'backward': # 1D equivalent: out[1:-1] = (f[1:-1] - f[:-2]) - np.subtract(f_arr[1:-1], f_arr[:-2], out=out[1:-1]) + arrmod.subtract(f_arr[1:-1], f_arr[:-2], out=out[1:-1]) # Boundaries if pad_mode == 'constant': diff --git a/odl/discr/discretization.py b/odl/discr/discretization.py index 14c5ec8515d..d714dab5639 100644 --- a/odl/discr/discretization.py +++ b/odl/discr/discretization.py @@ -331,16 +331,18 @@ def copy(self): """Create an identical (deep) copy of this element.""" return self.space.element(self.tensor.copy()) - def asarray(self, out=None): - """Extract the data of this array as a numpy array. + def asarray(self, out=None, impl='numpy'): + """Extract the data of this element as an array. Parameters ---------- - out : `numpy.ndarray`, optional - Array in which the result should be written in-place. - Has to be contiguous and of the correct dtype. + out : ndarray, optional + Array into which the result should be written. Must be contiguous + and of the correct dtype. + impl : str, optional + Array backend for the output, used when ``out`` is not given. """ - return self.tensor.asarray(out=out) + return self.tensor.asarray(out=out, impl=impl) def astype(self, dtype): """Return a copy of this element with new ``dtype``. diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index 1f6f74624e3..f93fec8acac 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -792,6 +792,11 @@ def element(self, inp=None, order=None): if 0 in arr.strides: arr = arr.copy() + # Make sure the shape is ok + if arr.shape != self.shape: + raise ValueError('`inp` must have shape {}, got shape {}' + ''.format(self.shape, arr.shape)) + return self.element_type(self, arr) def zero(self): @@ -1178,19 +1183,23 @@ def device(self): """The GPU device on which this tensor lies.""" return self.space.device - def asarray(self, out=None): - """Extract the data of this element as a `numpy.ndarray`. + def asarray(self, out=None, impl='numpy'): + """Extract the data of this element as an ndarray. + + This method is invoked when calling `numpy.asarray` on this tensor. Parameters ---------- - out : `numpy.ndarray`, optional - Array to which the result should be written. - Has to be contiguous and of the correct data type. + out : ndarray, optional + Array into which the result should be written. Must be contiguous + and of the correct dtype. + impl : str, optional + Array backend for the output, used when ``out`` is not given. Returns ------- - asarray : `numpy.ndarray` - Numpy array of the same `dtype` and `shape` this tensor. + asarray : ndarray + Array of the same `dtype` and `shape` this tensor. If ``out`` was given, the returned object is a reference to it. Examples @@ -1222,21 +1231,36 @@ def asarray(self, out=None): >>> result is out True """ + impl, impl_in = str(impl).lower(), impl if out is None: - return cupy.asnumpy(self.data) + if impl == 'numpy': + return cupy.asnumpy(self.data) + elif impl == 'cupy': + return self.data + else: + raise ValueError('`impl` {!r} not understood'.format(impl_in)) else: - if not isinstance(out, np.ndarray): - raise TypeError('`out` must be a `numpy.ndarray`, got type ' - '{}'.format(type(out))) + if not (out.flags.c_contiguous or out.flags.f_contiguous): + raise ValueError('`out` must be contiguous') if out.shape != self.shape: raise ValueError('`out` must have shape {}, got shape {}' ''.format(self.shape, out.shape)) if out.dtype != self.dtype: raise ValueError('`out` must have dtype {}, got dtype {}' ''.format(self.dtype, out.dtype)) - self.data.data.copy_to_host( - out.ctypes.data_as(np.ctypeslib.ctypes.c_void_p), - self.size * self.itemsize) + + if isinstance(out, np.ndarray): + if (self.data.flags.c_contiguous or + self.data.flags.f_contiguous): + # Use efficient copy for contiguous memory + self.data.data.copy_to_host( + out.ctypes.data_as(np.ctypeslib.ctypes.c_void_p), + self.size * self.itemsize) + else: + out[:] = cupy.asnumpy(self.data) + else: + out[:] = self.data + return out @property @@ -1473,7 +1497,7 @@ def __setitem__(self, indices, values): """ if isinstance(values, CupyTensor): self.data[indices] = values.data - elif np.isscalar(values): + elif isinstance(values, cupy.ndarray) or np.isscalar(values): self.data[indices] = values else: values = cupy.array(values, dtype=self.dtype, copy=False) @@ -2127,53 +2151,6 @@ def eval_at_via_npy(*inputs, **kwargs): return result - @property - def ufuncs(self): - """Access to NumPy style ufuncs. - - Examples - -------- - >>> r2 = odl.rn(2, impl='cupy') - >>> x = r2.element([1, -2]) - >>> x.ufuncs.absolute() - rn(2, impl='cupy').element([ 1., 2.]) - - These functions can also be used with broadcasting or - array-like input: - - >>> x.ufuncs.add(3) - rn(2, impl='cupy').element([ 4., 1.]) - >>> x.ufuncs.subtract([3, 3]) - rn(2, impl='cupy').element([-2., -5.]) - - There is also support for various reductions - (sum, prod, amin, amax): - - >>> x.ufuncs.sum() - -1.0 - >>> x.ufuncs.prod() - -2.0 - - They also support an out parameter - - >>> y = r2.element([3, 4]) - >>> out = r2.element() - >>> result = x.ufuncs.add(y, out=out) - >>> result - rn(2, impl='cupy').element([ 4., 2.]) - >>> result is out - True - - Notes - ----- - Those ufuncs which are implemented natively on the GPU incur no - significant overhead. However, for missing functions, a fallback - Numpy implementation is used which causes significant overhead - due to data copies between host and device. - """ - # TODO: Test with some native ufuncs, then remove this attribute - return super(CupyTensor, self).ufuncs - @property def real(self): """Real part of this tensor. @@ -2199,7 +2176,13 @@ def real(self, newreal): newreal : `array-like` or scalar The new real part for this tensor. """ - self.data.real[:] = newreal + # Assignment with array-likes does not work directly, see + # https://github.com/cupy/cupy/issues/593 + if np.isscalar(newreal): + self.data.real = newreal + else: + # Avoid compile errors for complex right-hand sides + self.data.real = cupy.asarray(newreal).real @property def imag(self): @@ -2226,7 +2209,13 @@ def imag(self, newimag): newimag : `array-like` or scalar The new imaginary part for this tensor. """ - self.data.imag[:] = newimag + # Assignment with array-likes does not work directly, see + # https://github.com/cupy/cupy/issues/593 + if np.isscalar(newimag): + self.data.imag = newimag + else: + # Avoid compile errors for complex right-hand sides + self.data.imag = cupy.asarray(newimag).real def conj(self, out=None): """Complex conjugate of this tensor. diff --git a/odl/space/npy_tensors.py b/odl/space/npy_tensors.py index 8a80bcc4b75..3edf978e77e 100644 --- a/odl/space/npy_tensors.py +++ b/odl/space/npy_tensors.py @@ -864,24 +864,24 @@ def data(self): """The `numpy.ndarray` representing the data of ``self``.""" return self.__data - def asarray(self, out=None): - """Extract the data of this array as a ``numpy.ndarray``. + def asarray(self, out=None, impl='numpy'): + """Extract the data of this array as an ndarray. - This method is invoked when calling `numpy.asarray` on this - tensor. + This method is invoked when calling `numpy.asarray` on this tensor. Parameters ---------- - out : `numpy.ndarray`, optional - Array in which the result should be written in-place. - Has to be contiguous and of the correct dtype. + out : ndarray, optional + Array into which the result should be written. Must be contiguous + and of the correct dtype. + impl : str, optional + Array backend for the output, used when ``out`` is not given. Returns ------- - asarray : `numpy.ndarray` - Numpy array with the same data type as ``self``. If - ``out`` was given, the returned object is a reference - to it. + asarray : ndarray + Array with the same data type as ``self``. If ``out`` was given, + the returned object is a reference to it. Examples -------- @@ -902,10 +902,45 @@ def asarray(self, out=None): array([[ 1., 1., 1.], [ 1., 1., 1.]]) """ + from odl.space.cupy_tensors import cupy, CUPY_AVAILABLE + import ctypes + + impl, impl_in = str(impl).lower(), impl if out is None: - return self.data + if impl == 'numpy': + return self.data + elif impl == 'cupy': + if CUPY_AVAILABLE: + return cupy.array(self.data) + else: + raise ValueError("`impl` 'cupy' not available") + else: + raise ValueError('`impl` {!r} not understood'.format(impl_in)) + else: - out[:] = self.data + if not (out.flags.c_contiguous or out.flags.f_contiguous): + raise ValueError('`out` must be contiguous') + if out.shape != self.shape: + raise ValueError('`out` must have shape {}, got shape {}' + ''.format(self.shape, out.shape)) + if out.dtype != self.dtype: + raise ValueError('`out` must have dtype {}, got dtype {}' + ''.format(self.dtype, out.dtype)) + + if CUPY_AVAILABLE and isinstance(out, cupy.ndarray): + # Use efficient copy by ensuring contiguous memory + if self.data.flags.contiguous: + self_contig_arr = self.data + else: + self_contig_arr = np.ascontiguousarray(self.data) + + out.data.copy_from_host( + self_contig_arr.ctypes.data_as(ctypes.c_void_p), + self.size * self.itemsize) + + else: + out[:] = self.data + return out def astype(self, dtype): diff --git a/odl/test/space/tensors_test.py b/odl/test/space/tensors_test.py index 888ca2dc34c..0ffbb680170 100644 --- a/odl/test/space/tensors_test.py +++ b/odl/test/space/tensors_test.py @@ -31,7 +31,7 @@ CUPY_AVAILABLE, cupy) from odl.util import array_module, array_cls, as_numpy from odl.util.testutils import ( - all_almost_equal, all_equal, simple_fixture, + all_almost_equal, all_equal, simple_fixture, skip_if_no_cupy, noise_array, noise_element, noise_elements, xfail_if) from odl.util.ufuncs import UFUNCS @@ -58,7 +58,7 @@ def _data_ptr(array): def _pos_array(space): """Create an array with positive real entries for ``space``.""" - return array_module(space.impl).abs(noise_array(space)) + 0.1 + return array_module(space.impl).asarray(abs(noise_array(space)) + 0.1) def _weighting_cls(impl, kind): @@ -331,7 +331,8 @@ def test_element(tspace, odl_elem_order): assert elem.data.flags[order + '_CONTIGUOUS'] # From array (C order) - arr_c = array_module(tspace.impl).ascontiguousarray(noise_array(tspace)) + arr_c = array_module(tspace.impl).asarray( + np.ascontiguousarray(noise_array(tspace))) elem = tspace.element(arr_c, order=order) assert all_equal(elem, arr_c) assert elem.shape == elem.data.shape @@ -346,7 +347,8 @@ def test_element(tspace, odl_elem_order): assert elem.data.flags[order + '_CONTIGUOUS'] # From array (F order) - arr_f = array_module(tspace.impl).asfortranarray(noise_array(tspace)) + arr_f = array_module(tspace.impl).asarray( + np.asfortranarray(noise_array(tspace))) elem = tspace.element(arr_f, order=order) assert all_equal(elem, arr_f) assert elem.shape == elem.data.shape @@ -929,7 +931,79 @@ def test_element_setitem_bool_array(odl_tspace_impl): assert all_equal(x, x_arr) +<<<<<<< e3e364b7a968160309e60aa4c2959d48fa2b6d7b def test_transpose(odl_tspace_impl): +======= +@skip_if_no_cupy +def test_asarray_numpy_to_cupy(floating_dtype): + """Test x.asarray with numpy x and cupy impl and out.""" + space = odl.tensor_space((2, 3), dtype=floating_dtype) + + with xfail_if(floating_dtype in ('float128', 'complex256'), + reason='quad precision types not available in cupy'): + # Make new array, contiguous + x = space.one() + x_cpy = x.asarray(impl='cupy') + assert isinstance(x_cpy, cupy.ndarray) + assert all_equal(x_cpy, x) + + # Write to existing, contiguous + out_cpy = cupy.empty((2, 3), dtype=floating_dtype) + x_cpy = x.asarray(out=out_cpy) + assert x_cpy is out_cpy + assert all_equal(out_cpy, x) + + # Make new array, discontiguous + arr = np.arange(12).astype(floating_dtype).reshape((2, 6))[:, ::2] + x = space.element(arr) + assert not (x.data.flags.c_contiguous or x.data.flags.f_contiguous) + x_cpy = x.asarray(impl='cupy') + assert isinstance(x_cpy, cupy.ndarray) + assert all_equal(x_cpy, x) + + # Write to existing, contiguous + out_cpy = cupy.empty((2, 3), dtype=floating_dtype) + x_cpy = x.asarray(out=out_cpy) + assert x_cpy is out_cpy + assert all_equal(out_cpy, x) + + +@skip_if_no_cupy +def test_asarray_cupy_to_numpy(floating_dtype): + """Test x.asarray with cupy x and numpy impl and out.""" + with xfail_if(floating_dtype in ('float128', 'complex256'), + reason='quad precision types not available in cupy'): + space = odl.tensor_space((2, 3), dtype=floating_dtype, impl='cupy') + + # Make new array, contiguous + x = space.one() + x_npy = x.asarray(impl='numpy') + assert isinstance(x_npy, np.ndarray) + assert all_equal(x_npy, x) + + # Write to existing, contiguous + out_npy = np.empty((2, 3), dtype=floating_dtype) + x_npy = x.asarray(out=out_npy) + assert x_npy is out_npy + assert all_equal(out_npy, x) + + # Make new array, discontiguous + arr = cupy.arange(12).astype(floating_dtype).reshape((2, 6))[:, ::2] + x = space.element(arr) + assert not (x.data.flags.c_contiguous or x.data.flags.f_contiguous) + x_npy = x.asarray(impl='numpy') + assert isinstance(x_npy, np.ndarray) + assert all_equal(x_npy, x) + + # Write to existing, contiguous + out_npy = np.empty((2, 3), dtype=floating_dtype) + x_npy = x.asarray(out=out_npy) + assert x_npy is out_npy + assert all_equal(out_npy, x) + + +def test_transpose(tspace_impl): +>>>>>>> ENH: add impl to asarray, tests, and fix real and imag for cupy """Test the .T property of tensors against plain inner product.""" impl = odl_tspace_impl spaces = [odl.rn((3, 4), impl=impl)] diff --git a/odl/util/testutils.py b/odl/util/testutils.py index 10d9d97f44f..e24288f7254 100644 --- a/odl/util/testutils.py +++ b/odl/util/testutils.py @@ -19,7 +19,6 @@ from odl.util.utility import run_from_ipython, is_string, none_context - __all__ = ( 'all_equal', 'all_almost_equal', 'dtype_ndigits', 'dtype_tol', 'never_skip', 'skip_if_no_stir', 'skip_if_no_pywavelets', @@ -28,7 +27,6 @@ 'ProgressRange', 'test', 'run_doctests', 'test_file' ) - def _ndigits(a, b, default=None): """Return number of expected correct digits comparing ``a`` and ``b``. @@ -225,6 +223,7 @@ def _pass(function): return function never_skip = _pass + skip_if_no_cupy = _pass skip_if_no_stir = _pass skip_if_no_pywavelets = _pass skip_if_no_pyfftw = _pass @@ -242,6 +241,11 @@ def _pass(function): reason='STIR not available' ) + skip_if_no_cupy = pytest.mark.skipif( + "not odl.space.cupy_tensors.CUPY_AVAILABLE", + reason='CuPy not available' + ) + skip_if_no_pywavelets = pytest.mark.skipif( "not odl.trafos.PYWT_AVAILABLE", reason='PyWavelets not available' From f4f376c10b70312065fadd73bdda22baaaccd71e Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Thu, 30 Nov 2017 11:20:15 +0100 Subject: [PATCH 29/38] WIP: add asarray helper, fix diff_ops for cupy --- odl/discr/diff_ops.py | 29 ++++----- odl/test/discr/diff_ops_test.py | 2 +- odl/test/space/tensors_test.py | 6 +- odl/util/utility.py | 108 +++++++++++++++++++++++++++++--- 4 files changed, 113 insertions(+), 32 deletions(-) diff --git a/odl/discr/diff_ops.py b/odl/discr/diff_ops.py index ba83586bf6b..a0e0b9aa4b2 100644 --- a/odl/discr/diff_ops.py +++ b/odl/discr/diff_ops.py @@ -14,7 +14,8 @@ from odl.discr.lp_discr import DiscreteLp from odl.operator.tensor_ops import PointwiseTensorFieldOperator from odl.space import ProductSpace -from odl.util import writable_array, signature_string, indent, array_module +from odl.util import ( + writable_array, asarray, array_module, signature_string, indent) __all__ = ('PartialDerivative', 'Gradient', 'Divergence', 'Laplacian') @@ -190,8 +191,7 @@ def __repr__(self): def __str__(self): """Return ``str(self)``.""" - dom_ran_str = '\n-->\n'.join([repr(self.domain), repr(self.range)]) - return '{}:\n{}'.format(self.__class__.__name__, indent(dom_ran_str)) + return repr(self) class Gradient(PointwiseTensorFieldOperator): @@ -347,14 +347,13 @@ def _call(self, x, out=None): if out is None: out = self.range.element() - x_arr = x.asarray() ndim = self.domain.ndim dx = self.domain.cell_sides impl = self.domain.impl for axis in range(ndim): with writable_array(out[axis], impl=impl) as out_arr: - finite_diff(x_arr, axis=axis, dx=dx[axis], impl=impl, + finite_diff(x, axis=axis, dx=dx[axis], impl=impl, method=self.method, pad_mode=self.pad_mode, pad_const=self.pad_const, out=out_arr) return out @@ -414,8 +413,7 @@ def __repr__(self): def __str__(self): """Return ``str(self)``.""" - dom_ran_str = '\n-->\n'.join([repr(self.domain), repr(self.range)]) - return '{}:\n{}'.format(self.__class__.__name__, indent(dom_ran_str)) + return repr(self) class Divergence(PointwiseTensorFieldOperator): @@ -625,8 +623,7 @@ def __repr__(self): def __str__(self): """Return ``str(self)``.""" - dom_ran_str = '\n-->\n'.join([repr(self.domain), repr(self.range)]) - return '{}:\n{}'.format(self.__class__.__name__, indent(dom_ran_str)) + return repr(self) class Laplacian(PointwiseTensorFieldOperator): @@ -719,23 +716,20 @@ def _call(self, x, out=None): ndim = self.domain.ndim dx = self.domain.cell_sides impl = self.domain.impl - - x_arr = x.asarray(impl=impl) - out_arr = out.asarray(impl=impl) tmp = array_module(impl).empty( out.shape, out.dtype, order=out.space.default_order) with writable_array(out, impl=impl) as out_arr: for axis in range(ndim): # TODO: this can be optimized - finite_diff(x_arr, axis=axis, dx=dx[axis] ** 2, impl=impl, + finite_diff(x, axis=axis, dx=dx[axis] ** 2, impl=impl, method='forward', pad_mode=self.pad_mode, pad_const=self.pad_const, out=tmp) out_arr += tmp - finite_diff(x_arr, axis=axis, dx=dx[axis] ** 2, impl=impl, + finite_diff(x, axis=axis, dx=dx[axis] ** 2, impl=impl, method='backward', pad_mode=self.pad_mode, pad_const=self.pad_const, out=tmp) @@ -785,8 +779,7 @@ def __repr__(self): def __str__(self): """Return ``str(self)``.""" - dom_ran_str = '\n-->\n'.join([repr(self.domain), repr(self.range)]) - return '{}:\n{}'.format(self.__class__.__name__, indent(dom_ran_str)) + return repr(self) def finite_diff(f, axis, dx=1.0, method='forward', out=None, impl='numpy', @@ -891,8 +884,7 @@ def finite_diff(f, axis, dx=1.0, method='forward', out=None, impl='numpy', >>> out is finite_diff(f, axis=0, out=out) True """ - arrmod = array_module(impl) - f_arr = arrmod.asarray(f) + f_arr = asarray(f, impl=impl) ndim = f_arr.ndim if f_arr.shape[axis] < 2: @@ -921,6 +913,7 @@ def finite_diff(f, axis, dx=1.0, method='forward', out=None, impl='numpy', pad_const = kwargs.pop('pad_const', 0) pad_const = f.dtype.type(pad_const) + arrmod = array_module(impl) if out is None: out = arrmod.empty_like(f_arr) else: diff --git a/odl/test/discr/diff_ops_test.py b/odl/test/discr/diff_ops_test.py index 90d0a910ccc..ff8a759afd2 100644 --- a/odl/test/discr/diff_ops_test.py +++ b/odl/test/discr/diff_ops_test.py @@ -28,7 +28,7 @@ 'order0', 'order1', 'order2']) -@pytest.fixture(scope="module", params=[1, 2, 3], ids=['1d', '2d', '3d']) +@pytest.fixture(scope="module", params=[1, 2, 3], ids=[' 1d ', ' 2d ', ' 3d ']) def space(request, odl_tspace_impl): impl = odl_tspace_impl ndim = request.param diff --git a/odl/test/space/tensors_test.py b/odl/test/space/tensors_test.py index 0ffbb680170..71fb8ed8963 100644 --- a/odl/test/space/tensors_test.py +++ b/odl/test/space/tensors_test.py @@ -931,9 +931,6 @@ def test_element_setitem_bool_array(odl_tspace_impl): assert all_equal(x, x_arr) -<<<<<<< e3e364b7a968160309e60aa4c2959d48fa2b6d7b -def test_transpose(odl_tspace_impl): -======= @skip_if_no_cupy def test_asarray_numpy_to_cupy(floating_dtype): """Test x.asarray with numpy x and cupy impl and out.""" @@ -1002,8 +999,7 @@ def test_asarray_cupy_to_numpy(floating_dtype): assert all_equal(out_npy, x) -def test_transpose(tspace_impl): ->>>>>>> ENH: add impl to asarray, tests, and fix real and imag for cupy +def test_transpose(odl_tspace_impl): """Test the .T property of tensors against plain inner product.""" impl = odl_tspace_impl spaces = [odl.rn((3, 4), impl=impl)] diff --git a/odl/util/utility.py b/odl/util/utility.py index b91733c4fdb..d06dcb278d3 100644 --- a/odl/util/utility.py +++ b/odl/util/utility.py @@ -31,14 +31,34 @@ 'cache_arguments', 'unique', 'REPR_PRECISION') -__all__ = ('array_str', 'dtype_str', 'dtype_repr', 'npy_printoptions', - 'signature_string', 'indent', - 'is_numeric_dtype', 'is_int_dtype', 'is_floating_dtype', - 'is_real_dtype', 'is_real_floating_dtype', - 'is_complex_floating_dtype', 'real_dtype', 'complex_dtype', - 'is_string', 'conj_exponent', 'writable_array', 'none_context', - 'array_module', 'array_cls', 'as_numpy', - 'run_from_ipython', 'NumpyRandomSeed', 'cache_arguments', 'unique') +__all__ = ( + 'NumpyRandomSeed', + 'array_cls', + 'array_module', + 'array_str', + 'as_numpy', + 'asarray', + 'cache_arguments', + 'complex_dtype', + 'conj_exponent', + 'dtype_repr', + 'dtype_str', + 'indent', + 'is_complex_floating_dtype', + 'is_floating_dtype', + 'is_int_dtype', + 'is_numeric_dtype', + 'is_real_dtype', + 'is_real_floating_dtype', + 'is_string', + 'none_context', + 'npy_printoptions', + 'real_dtype', + 'run_from_ipython', + 'signature_string', + 'unique', + 'writable_array' + ) REPR_PRECISION = 4 # For printing scalars and array entries TYPE_MAP_R2C = {np.dtype(dtype): np.result_type(dtype, 1j) @@ -656,6 +676,78 @@ def ip_wrapper(x, out, **kwargs): return decorator +def asarray(obj, dtype=None, impl='numpy'): + """Convert ``obj`` to an ``impl`` type array as fast as possible. + + Parameters + ---------- + obj : array_like + Object to be converted to an array. + dtype : data-type, optional + Desired data type of the array. + impl : str, optional + Array backend used to create the array. + + Returns + ------- + array : ndarray + Array with data type ``dtype`` created from ``obj`` using ``impl`` + as backend. + + Examples + -------- + >>> a = asarray([1, 2]) + >>> a + array([1, 2]) + >>> type(a) + numpy.ndarray + >>> a = asarray(odl.rn(3).one()) + >>> a + array([ 1., 1., 1.]) + >>> type(a) + numpy.ndarray + + Notes + ----- + For ODL tensors, there are specific implementation of the conversion + to different types of arrays that are not necessarily invoked by + the ``asarray`` methods of array libraries. + While ``numpy.asarray`` uses the ``__array__`` method to "ask" the + object for an array, there is no such dedicated interface for other + implementations. As a result, an intermediate Numpy array is typically + created, which, for instance, for ``impl='cupy'`` leads to transfers + GPU->CPU->GPU. + This function bypasses this detour and uses the optimized + ``Tensor.asarray()`` method. + """ + from odl.space.cupy_tensors import cupy, CUPY_AVAILABLE + impl, impl_in = str(impl).lower(), impl + + if impl == 'numpy': + if CUPY_AVAILABLE and isinstance(obj, cupy.ndarray): + # __array__ of cupy.ndarray does not return a Numpy array, see + # https://github.com/cupy/cupy/issues/589 + obj = cupy.asnumpy(obj) + return np.asarray(obj, dtype=dtype) + + elif impl == 'cupy': + if CUPY_AVAILABLE: + try: + arr = obj.asarray(impl='cupy') + except AttributeError: + arr = cupy.asarray(obj, dtype=dtype) + else: + arr = arr.astype(dtype, copy=False) + + return arr + + else: + raise ValueError("`impl` 'cupy' not available") + + else: + raise ValueError('`impl` {!r} not understood'.format(impl_in)) + + class writable_array(object): """Context manager that casts obj to a `numpy.array` and saves changes.""" From d0fc8e2adcd82178d793424b41e46a1671687146 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Fri, 1 Dec 2017 00:45:39 +0100 Subject: [PATCH 30/38] MAINT: fix utils --- odl/util/testutils.py | 32 ++++++++++++++++++++------------ odl/util/utility.py | 4 ++-- 2 files changed, 22 insertions(+), 14 deletions(-) diff --git a/odl/util/testutils.py b/odl/util/testutils.py index e24288f7254..a86204fe139 100644 --- a/odl/util/testutils.py +++ b/odl/util/testutils.py @@ -17,7 +17,7 @@ import warnings from time import time -from odl.util.utility import run_from_ipython, is_string, none_context +from odl.util.utility import run_from_ipython, is_string, none_context, asarray __all__ = ( 'all_equal', 'all_almost_equal', 'dtype_ndigits', 'dtype_tol', @@ -362,17 +362,19 @@ def noise_array(space): typical to the space. """ from odl.space.pspace import ProductSpace - from odl.space.cupy_tensors import cupy if isinstance(space, ProductSpace): arr_list = [noise_array(spc_i) for spc_i in space] if space.is_power_space: arr = np.empty((len(arr_list),) + arr_list[0].shape, - dtype=space.dtype) - for i in range(len(arr)): - arr[i] = arr_list[i] + dtype=space[0].dtype) else: - return tuple(arr_list) + arr = np.empty((len(arr_list),) + arr_list[0].shape, + dtype=object) + for i in range(len(arr)): + arr[i] = arr_list[i] + + return arr else: if space.dtype == bool: @@ -483,15 +485,21 @@ def noise_elements(space, n=1): noise_array noise_element """ - npy_arrs = [noise_array(space) for _ in range(n)] + from odl.space.pspace import ProductSpace + + if isinstance(space, ProductSpace) and not space.is_power_space: + raise ValueError('`space` cannot be a non-power product space') + + if isinstance(space, ProductSpace): + impl = space[0].impl + else: + impl = space.impl - # Make space elements from Numpy arrays - elems = tuple(space.element(arr) for arr in npy_arrs) - # Make copies of the arrays - arrs = tuple(elem.data.copy() for elem in elems) + arrs = tuple(asarray(noise_array(space), impl=impl) for _ in range(n)) + elems = tuple(space.element(arr) for arr in arrs) if n == 1: - return tuple(arrs + elems) + return arrs + elems else: return arrs, elems diff --git a/odl/util/utility.py b/odl/util/utility.py index d06dcb278d3..56ea21b8728 100644 --- a/odl/util/utility.py +++ b/odl/util/utility.py @@ -700,12 +700,12 @@ def asarray(obj, dtype=None, impl='numpy'): >>> a array([1, 2]) >>> type(a) - numpy.ndarray + >>> a = asarray(odl.rn(3).one()) >>> a array([ 1., 1., 1.]) >>> type(a) - numpy.ndarray + Notes ----- From c632802bc66cd5da795b7fad5a49713c2e7fa7ab Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Fri, 1 Dec 2017 00:46:11 +0100 Subject: [PATCH 31/38] MAINT: exclude unsupported ufunc operators --- odl/ufunc_ops/ufunc_ops.py | 45 ++++++++++++++++++++++++-------------- odl/util/ufuncs.py | 2 +- 2 files changed, 29 insertions(+), 18 deletions(-) diff --git a/odl/ufunc_ops/ufunc_ops.py b/odl/ufunc_ops/ufunc_ops.py index 2e66b711240..b2b541cd455 100644 --- a/odl/ufunc_ops/ufunc_ops.py +++ b/odl/ufunc_ops/ufunc_ops.py @@ -227,18 +227,19 @@ def __init__(self, space): if nin == 1: domain = space0 = space - dtypes = [space.dtype] - elif nin == len(space) == 2 and isinstance(space, ProductSpace): + dtypes_in = [space.dtype] + elif nin == 2: + if not (isinstance(space, ProductSpace) and len(space) == 2): + raise TypeError('`space` must be a `ProductSpace` of length ' + '{} for ufunc {!r}, got {!r}' + ''.format(nin, name, space)) domain = space space0 = space[0] - dtypes = [space[0].dtype, space[1].dtype] + dtypes_in = [space[0].dtype, space[1].dtype] else: - domain = ProductSpace(space, nin) - space0 = space - dtypes = [space.dtype, space.dtype] + raise RuntimeError('bad `nin` {}'.format(nin)) - dts_out = dtypes_out(name, dtypes) - print(dts_out) + dts_out = dtypes_out(name, dtypes_in) if nout == 1: range = space0.astype(dts_out[0]) @@ -251,7 +252,7 @@ def __init__(self, space): def _call(self, x, out=None): """Return ``self(x)``.""" - # TODO: use `__array_ufunc__` when implemented on `ProductSpace`, + # TODO: use `__array_ufunc__` when implemented on product spaces, # or try both if out is None: if nin == 1: @@ -288,8 +289,9 @@ def __repr__(self): result = getattr(vec.ufuncs, name)(vec2) if nout == 2: - result_space = ProductSpace(vec.space, 2) - result = repr(result_space.element(result)) + result_space = ProductSpace(result[0].space, result[1].space) + with np.errstate(all='ignore'): + result = repr(result_space.element(result)) examples_docstring = RAW_EXAMPLES_DOCSTRING.format(space=space, name=name, arg=arg, result=result) @@ -342,6 +344,7 @@ def __repr__(self): # Create example (also functions as doctest) + # TODO: remove this restriction if nin != 1: raise NotImplementedError('Currently not suppored') @@ -372,7 +375,7 @@ def __repr__(self): RAW_UFUNC_FACTORY_DOCSTRING = """{docstring} Notes ----- -This creates a `Operator`/`Functional` that applies a ufunc pointwise. +This creates an `Operator`or `Functional` that applies a ufunc pointwise. Examples -------- @@ -381,21 +384,28 @@ def __repr__(self): """ RAW_UFUNC_FACTORY_FUNCTIONAL_DOCSTRING = """ -Create functional with domain/range as real numbers: +Create functional with real numbers as domain/range: >>> func = odl.ufunc_ops.{name}() """ RAW_UFUNC_FACTORY_OPERATOR_DOCSTRING = """ -Create operator that acts pointwise on a `TensorSpace` +Create operator that acts pointwise on a `TensorSpace`: >>> space = odl.rn(3) >>> op = odl.ufunc_ops.{name}(space) """ +# Avoid tons of warnings on the console when testing +npy_err_old = np.geterr() +np.seterr(all='ignore') # Create an operator for each ufunc for name, nin, nout, docstring in UFUNCS: + if nin == 2: + # Currently not supported + continue + def indirection(name, docstring): # Indirection is needed since name should be saved but is changed # in the loop. @@ -413,11 +423,11 @@ def ufunc_factory(domain=RealNumbers()): globals()[name + '_op'] = ufunc_class_factory(name, nin, nout, docstring) - if not _is_integer_only_ufunc(name): + if _is_integer_only_ufunc(name): + operator_example = "" + else: operator_example = RAW_UFUNC_FACTORY_OPERATOR_DOCSTRING.format( name=name) - else: - operator_example = "" if not _is_integer_only_ufunc(name) and nin == 1 and nout == 1: globals()[name + '_func'] = ufunc_functional_factory( @@ -437,6 +447,7 @@ def ufunc_factory(domain=RealNumbers()): globals()[name] = ufunc_factory __all__ += (name,) +np.seterr(**npy_err_old) if __name__ == '__main__': from odl.util.testutils import run_doctests diff --git a/odl/util/ufuncs.py b/odl/util/ufuncs.py index 2545d17d252..2a6b50b92eb 100644 --- a/odl/util/ufuncs.py +++ b/odl/util/ufuncs.py @@ -41,7 +41,7 @@ 'conj', 'conjugate', 'copysign', 'cos', 'cosh', 'deg2rad', 'degrees', 'divide', 'equal', 'exp', 'exp2', 'expm1', 'fabs', 'floor', 'floor_divide', 'fmax', 'fmin', 'fmod', 'frexp', 'greater', 'greater_equal', 'hypot', - 'invert', 'isfinite', 'isinf', 'isnan', 'ldexp', 'left_shift', 'less', + 'invert', 'isfinite', 'isinf', 'isnan', 'left_shift', 'less', 'less_equal', 'log', 'log10', 'log1p', 'log2', 'logaddexp', 'logaddexp2', 'logical_and', 'logical_not', 'logical_or', 'logical_xor', 'maximum', 'minimum', 'mod', 'modf', 'multiply', 'negative', 'nextafter', From 7879a50cf0df00cac6d96f43374f8c7bdb1c7043 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Fri, 1 Dec 2017 00:46:32 +0100 Subject: [PATCH 32/38] MAINT: fix bad parens in test --- odl/test/space/pspace_test.py | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/odl/test/space/pspace_test.py b/odl/test/space/pspace_test.py index 75b4b17358d..c0ca2b48511 100644 --- a/odl/test/space/pspace_test.py +++ b/odl/test/space/pspace_test.py @@ -262,7 +262,8 @@ def test_metric(): dist = HxH.dist(w1, w2) expected_dist = np.sqrt( - sum(d ** 2) for d in (H.dist(v11, v21), H.dist(v12, v22))) + sum(d ** 2 for d in (H.dist(v11, v21), H.dist(v12, v22))) + ) assert dist == pytest.approx(expected_dist) # inf norm @@ -294,7 +295,8 @@ def test_norm(): norm = HxH.norm(w) expected_norm = np.sqrt( - sum(n ** 2) for n in (H.norm(v1), H.norm(v2))) + sum(n ** 2 for n in (H.norm(v1), H.norm(v2))) + ) assert norm == pytest.approx(expected_norm) # inf norm From 4d1183e98b8d68e1e5add13c27a23fea5f12905c Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Fri, 1 Dec 2017 00:47:20 +0100 Subject: [PATCH 33/38] BUG: fix indexing with cupy tensor --- odl/discr/discretization.py | 2 +- odl/discr/lp_discr.py | 23 ----------------------- odl/space/cupy_tensors.py | 3 +++ 3 files changed, 4 insertions(+), 24 deletions(-) diff --git a/odl/discr/discretization.py b/odl/discr/discretization.py index d714dab5639..84c8c860bfd 100644 --- a/odl/discr/discretization.py +++ b/odl/discr/discretization.py @@ -411,7 +411,7 @@ def __setitem__(self, indices, values): indices = indices.tensor if isinstance(values, type(self)): values = values.tensor - self.tensor.__setitem__(indices, values) + self.tensor[indices] = values def sampling(self, ufunc, **kwargs): """Sample a continuous function and assign to this element. diff --git a/odl/discr/lp_discr.py b/odl/discr/lp_discr.py index 54a849b2ff4..eaea13ef1c7 100644 --- a/odl/discr/lp_discr.py +++ b/odl/discr/lp_discr.py @@ -755,29 +755,6 @@ def conj(self, out=None): self.tensor.conj(out=out.tensor) return out - def __setitem__(self, indices, values): - """Set values of this element. - - Parameters - ---------- - indices : int or `slice` - The position(s) that should be set - values : scalar or `array-like` - Value(s) to be assigned. - If ``indices`` is an integer, ``values`` must be a scalar - value. - If ``indices`` is a slice, ``values`` must be - broadcastable to the size of the slice (same size, - shape ``(1,)`` or scalar). - For ``indices == slice(None)``, i.e. in the call - ``vec[:] = values``, a multi-dimensional array of correct - shape is allowed as ``values``. - """ - if values in self.space: - self.tensor[indices] = values.tensor - else: - super(DiscreteLpElement, self).__setitem__(indices, values) - def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): """Interface to Numpy's ufunc machinery. diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index f93fec8acac..4f046415628 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -1495,6 +1495,9 @@ def __setitem__(self, indices, values): >>> x[0] 4294967295 """ + if isinstance(indices, CupyTensor): + indices = indices.data + if isinstance(values, CupyTensor): self.data[indices] = values.data elif isinstance(values, cupy.ndarray) or np.isscalar(values): From 4ac503dcf07b213f2edca98900bc82c6bd4d2017 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Fri, 1 Dec 2017 00:47:43 +0100 Subject: [PATCH 34/38] MAINT: minor fix in diff_ops --- odl/discr/diff_ops.py | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/odl/discr/diff_ops.py b/odl/discr/diff_ops.py index a0e0b9aa4b2..e78c2b76dfa 100644 --- a/odl/discr/diff_ops.py +++ b/odl/discr/diff_ops.py @@ -140,7 +140,7 @@ def _call(self, x, out=None): impl = self.domain.impl with writable_array(out, impl=impl) as out_arr: - finite_diff(x.asarray(), axis=self.axis, dx=self.dx, impl=impl, + finite_diff(x, axis=self.axis, dx=self.dx, impl=impl, method=self.method, pad_mode=self.pad_mode, pad_const=self.pad_const, out=out_arr) return out @@ -492,10 +492,12 @@ def __init__(self, domain=None, range=None, method='forward', ... [2., 3., 4., 5., 6.]]) >>> f = div.domain.element([data, data]) >>> div_f = div(f) - >>> print(div_f) - [[ 2., 2., 2., 2., -3.], - [ 2., 2., 2., 2., -4.], - [ -1., -2., -3., -4., -12.]] + >>> div_f + uniform_discr([ 0., 0.], [ 3., 5.], (3, 5)).element( + [[ 2., 2., 2., 2., -3.], + [ 2., 2., 2., 2., -4.], + [ -1., -2., -3., -4., -12.]] + ) Verify adjoint: From 5c638d6d22d234cda56f69e71e91e67d0f6d3b92 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Fri, 1 Dec 2017 09:31:40 +0100 Subject: [PATCH 35/38] TST: mark cupy-only test as skip if cupy is not available --- odl/test/space/tensors_test.py | 1 + 1 file changed, 1 insertion(+) diff --git a/odl/test/space/tensors_test.py b/odl/test/space/tensors_test.py index 71fb8ed8963..8c8a1a6b1bb 100644 --- a/odl/test/space/tensors_test.py +++ b/odl/test/space/tensors_test.py @@ -1758,6 +1758,7 @@ def _check_result_is_out(result, out_seq): np.seterr(**npy_err_orig) +@skip_if_no_cupy def test_ufunc_cupy_force_native(): """Test the ``force_native`` flag for cupy based ufuncs.""" if not USE_ARRAY_UFUNCS_INTERFACE: From 3cdde9fbd8c3712cbac2268e6071896a56aa2cf6 Mon Sep 17 00:00:00 2001 From: Holger Kohr Date: Fri, 1 Dec 2017 10:27:58 +0100 Subject: [PATCH 36/38] MAINT: fix Py2-incompatible doctest in utility --- odl/util/utility.py | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/odl/util/utility.py b/odl/util/utility.py index 56ea21b8728..73e2fbe5137 100644 --- a/odl/util/utility.py +++ b/odl/util/utility.py @@ -699,13 +699,13 @@ def asarray(obj, dtype=None, impl='numpy'): >>> a = asarray([1, 2]) >>> a array([1, 2]) - >>> type(a) - + >>> isinstance(a, np.ndarray) + True >>> a = asarray(odl.rn(3).one()) >>> a array([ 1., 1., 1.]) - >>> type(a) - + >>> isinstance(a, np.ndarray) + True Notes ----- From 87f8401358f640a6f804b70ee78883211b2dcb2c Mon Sep 17 00:00:00 2001 From: Jonas Adler Date: Wed, 14 Feb 2018 11:17:49 +0100 Subject: [PATCH 37/38] MAINT: Minor improvements related to cupy --- odl/space/cupy_tensors.py | 102 ++++++++++++++++---------------- odl/space/npy_tensors.py | 17 ++++-- odl/test/discr/lp_discr_test.py | 17 +++--- odl/test/space/pspace_test.py | 11 ++-- odl/util/testutils.py | 11 +++- 5 files changed, 87 insertions(+), 71 deletions(-) diff --git a/odl/space/cupy_tensors.py b/odl/space/cupy_tensors.py index 4f046415628..6d534f63f6b 100644 --- a/odl/space/cupy_tensors.py +++ b/odl/space/cupy_tensors.py @@ -80,7 +80,7 @@ def _get_flat_inc(arr1, arr2=None): be applied, a ``ValueError`` is raised, triggering a fallback implementation. - For **1 array**, the conditions to be fulfilled are + For ``arr2 is None``, the conditions to be fulfilled are - the strides do not contain 0 and - the memory of the array has constant stride. @@ -93,9 +93,9 @@ def _get_flat_inc(arr1, arr2=None): - strided slices along the **fastest-varying axis** (``arr[..., ::2]`` for C-contiguous ``arr``). - For **2 arrays**, both arrays must + For ``arr2 is not None``, both arrays must - - fulfill the "1 array" conditions individually, + - fulfill the ``arr2 is None`` conditions individually, - have the same total size, and - the axis order of both must be the same in the sense that the same index array sorts the strides of both arrays in ascending order. @@ -298,17 +298,16 @@ def _get_scal(arr): return _fallback_scal try: - _cublas_scal = _cublas_func('scal', arr.dtype) + _cublas_scal_impl = _cublas_func('scal', arr.dtype) except (ValueError, AttributeError): return _fallback_scal - def scal(a, x): + def _cublas_scal(a, x): """Implement ``x <- a * x`` with constant ``a``.""" - return _cublas_scal( + return _cublas_scal_impl( x.data.device.cublas_handle, x.size, a, x.data.ptr, incx) - scal.__name__ = scal.__qualname__ = '_cublas_scal' - return scal + return _cublas_scal def _get_axpy(arr1, arr2): @@ -480,14 +479,14 @@ def __init__(self, shape, dtype='float64', device=None, **kwargs): Parameters ---------- - shape : sequence of non-negative ints - Number entries per dimension. + shape : positive int or sequence of positive ints + Number of entries per axis for elements in this space. + A single integer results in a space with rank 1, i.e., 1 axis. dtype : - Data type for each tuple entry. Can be provided in any - way the `numpy.dtype` function understands, e.g., - as built-in type, as one of NumPy's internal datatype - objects or as string. - See `available_dtypes` for the list of supported data types. + Data type of each element. Can be provided in any + way the `numpy.dtype` function understands, e.g. + as built-in type or as a string. + See `available_dtypes` for the list of supported scalar data types. device : int, optional ID of the GPU device where elements should be created. For ``None``, the default device is chosen, which usually @@ -512,7 +511,7 @@ def __init__(self, shape, dtype='float64', device=None, **kwargs): This option cannot be combined with ``dist``, ``norm`` or ``inner``. - Default: no weighting + Default: Each point has weight 1.0 (no weighting) exponent : positive float, optional Exponent of the norm. For values other than 2.0, no @@ -875,7 +874,8 @@ def __eq__(self, other): def __hash__(self): """Return ``hash(self)``.""" - return hash((super(CupyTensorSpace, self).__hash__(), self.device, + return hash((super(CupyTensorSpace, self).__hash__(), + self.device, self.weighting)) def _lincomb(self, a, x1, b, x2, out): @@ -1122,12 +1122,20 @@ def available_dtypes(): """Return the data types available for this space.""" dtypes = (np.sctypes['float'] + np.sctypes['complex'] + - np.sctypes['int'] - + np.sctypes['uint'] - + [bool]) - dtypes.remove(np.float128) - dtypes.remove(np.complex256) - return tuple(np.dtype(dtype) for dtype in dtypes) + np.sctypes['int'] + + np.sctypes['uint'] + + [bool]) + + # Convert to dtypes and remove duplicates + dtypes = tuple(set(np.dtype(dtype) for dtype in dtypes)) + + # Remove float128 and complex256 if they exist + if hasattr(np, 'float128'): + dtypes.remove(np.float128) + if hasattr(np, 'complex256'): + dtypes.remove(np.complex256) + + return dtypes @staticmethod def default_dtype(field=None): @@ -1267,6 +1275,12 @@ def asarray(self, out=None, impl='numpy'): def data_ptr(self): """Memory address of the data container as 64-bit integer. + Returns + ------- + data_ptr : int + The data pointer is technically of type ``uintptr_t`` and gives the + index in bytes of the first element of the data storage. + Examples -------- >>> r3 = odl.rn(3, impl='cupy') @@ -1290,11 +1304,6 @@ def __eq__(self, other): ``True`` if all entries of ``other`` are equal to this tensor's entries, ``False`` otherwise. - Notes - ----- - The element-by-element comparison is performed on the CPU, - i.e. it involves data transfer to host memory, which is slow. - Examples -------- >>> r3 = odl.rn(3, impl='cupy') @@ -1611,7 +1620,7 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): - ``np.add.reduce`` -> ``cupy.sum`` - ``np.add.accumulate`` -> ``cupy.cumsum`` - ``np.multiply.reduce`` -> ``cupy.prod`` - - ``np.multiply.reduce`` -> ``cupy.cumprod``. + - ``np.multiply.accumulate`` -> ``cupy.cumprod``. **All other such methods will run Numpy code and be slow**! @@ -1853,14 +1862,10 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): # Manually implemented cases if ( (ufunc in (np.add, np.multiply) and - method in ('reduce', 'accumulate') - ) or + method in ('reduce', 'accumulate')) or (ufunc in (np.minimum, np.maximum) and - method == 'reduce' - ) or - (ufunc == np.add and - method == 'at' - ) + method == 'reduce') or + (ufunc == np.add and method == 'at') ): use_native = native_ufunc is not None # should be True else: @@ -1902,19 +1907,6 @@ def __array_ufunc__(self, ufunc, method, *inputs, **kwargs): elif isinstance(inputs[i], CupyTensor): inputs[i] = cupy.asnumpy(inputs[i].data) - # For debugging - if use_native: - assert all(isinstance(i, cupy.ndarray) or - np.isscalar(i) or - i is None - for i in inputs) - assert all(isinstance(o, cupy.ndarray) or o is None - for o in (out1, out2)) - - elif not use_native and method != 'at': - assert not any(isinstance(i, cupy.ndarray) for i in inputs) - # `out` handled below, we don't want to mess with it here - # --- For later --- # # Wrap the result in an appropriate space, propagating weighting @@ -2290,6 +2282,14 @@ def _weighting(weights, exponent): if CUPY_AVAILABLE: + dot = cupy.ReductionKernel(in_params='T x, T y', + out_params='T res', + map_expr='x * y', + reduce_expr='a + b', + post_map_expr='res = a', + identity='0', + name='dot') + dotw = cupy.ReductionKernel(in_params='T x, T y, W w', out_params='T res', map_expr='x * y * w', @@ -2396,7 +2396,7 @@ def _weighting(weights, exponent): dist2 = cupy.ReductionKernel(in_params='T x, T y', out_params='R res', - map_expr='abs(x - y) * abs(x - y)', + map_expr='abs((x - y) * (x - y))', reduce_expr='a + b', post_map_expr='res = sqrt(a)', identity='0', @@ -2404,7 +2404,7 @@ def _weighting(weights, exponent): dist2w = cupy.ReductionKernel(in_params='T x, T y, W w', out_params='R res', - map_expr='abs(x - y) * abs(x - y) * w', + map_expr='abs((x - y) * (x - y)) * w', reduce_expr='a + b', post_map_expr='res = sqrt(a)', identity='0', diff --git a/odl/space/npy_tensors.py b/odl/space/npy_tensors.py index 3edf978e77e..c2e9158a59a 100644 --- a/odl/space/npy_tensors.py +++ b/odl/space/npy_tensors.py @@ -73,19 +73,19 @@ class NumpyTensorSpace(TensorSpace): .. _Wikipedia article on tensors: https://en.wikipedia.org/wiki/Tensor """ - def __init__(self, shape, dtype=None, **kwargs): + def __init__(self, shape, dtype='float64', **kwargs): """Initialize a new instance. Parameters ---------- shape : positive int or sequence of positive ints - Number of entries per axis for elements in this space. A - single integer results in a space with rank 1, i.e., 1 axis. + Number of entries per axis for elements in this space. + A single integer results in a space with rank 1, i.e., 1 axis. dtype : Data type of each element. Can be provided in any way the `numpy.dtype` function understands, e.g. - as built-in type or as a string. For ``None``, - the `default_dtype` of this space (``float64``) is used. + as built-in type or as a string. + See `available_dtypes` for the list of supported scalar data types. exponent : positive float, optional Exponent of the norm. For values other than 2.0, no inner product is defined. @@ -902,6 +902,7 @@ def asarray(self, out=None, impl='numpy'): array([[ 1., 1., 1.], [ 1., 1., 1.]]) """ + # Import cupy if it exists (else None) from odl.space.cupy_tensors import cupy, CUPY_AVAILABLE import ctypes @@ -965,6 +966,12 @@ def astype(self, dtype): def data_ptr(self): """Memory address of the data container as 64-bit integer. + Returns + ------- + data_ptr : int + The data pointer is technically of type ``uintptr_t`` and gives the + index in bytes of the first element of the data storage. + Examples -------- >>> import ctypes diff --git a/odl/test/discr/lp_discr_test.py b/odl/test/discr/lp_discr_test.py index 3737b116f41..280ccb3aa6c 100644 --- a/odl/test/discr/lp_discr_test.py +++ b/odl/test/discr/lp_discr_test.py @@ -744,15 +744,14 @@ def test_cell_volume(): def test_astype(tspace_impl): - - rdiscr = odl.uniform_discr([0, 0], [1, 1], [2, 2], dtype='float64', - impl=tspace_impl) - cdiscr = odl.uniform_discr([0, 0], [1, 1], [2, 2], dtype='complex128', - impl=tspace_impl) - rdiscr_s = odl.uniform_discr([0, 0], [1, 1], [2, 2], dtype='float32', - impl=tspace_impl) - cdiscr_s = odl.uniform_discr([0, 0], [1, 1], [2, 2], dtype='complex64', - impl=tspace_impl) + rdiscr = odl.uniform_discr([0, 0], [1, 1], [2, 2], + dtype='float64', impl=tspace_impl) + cdiscr = odl.uniform_discr([0, 0], [1, 1], [2, 2], + dtype='complex128', impl=tspace_impl) + rdiscr_s = odl.uniform_discr([0, 0], [1, 1], [2, 2], + dtype='float32', impl=tspace_impl) + cdiscr_s = odl.uniform_discr([0, 0], [1, 1], [2, 2], + dtype='complex64', impl=tspace_impl) # Real assert rdiscr.astype('float32') == rdiscr_s diff --git a/odl/test/space/pspace_test.py b/odl/test/space/pspace_test.py index c0ca2b48511..d6bce1858c4 100644 --- a/odl/test/space/pspace_test.py +++ b/odl/test/space/pspace_test.py @@ -252,7 +252,8 @@ def test_metric(): w2 = HxH.element([v21, v22]) dist = HxH.dist(w1, w2) - expected_dist = sum([H.dist(v11, v21), H.dist(v12, v22)]) + expected_dist = np.linalg.norm([H.dist(v11, v21), H.dist(v12, v22)], + ord=1.0) assert dist == pytest.approx(expected_dist) # 2-norm @@ -261,9 +262,8 @@ def test_metric(): w2 = HxH.element([v21, v22]) dist = HxH.dist(w1, w2) - expected_dist = np.sqrt( - sum(d ** 2 for d in (H.dist(v11, v21), H.dist(v12, v22))) - ) + expected_dist = np.linalg.norm([H.dist(v11, v21), H.dist(v12, v22)], + ord=2.0) assert dist == pytest.approx(expected_dist) # inf norm @@ -272,7 +272,8 @@ def test_metric(): w2 = HxH.element([v21, v22]) dist = HxH.dist(w1, w2) - expected_dist = max(H.dist(v11, v21), H.dist(v12, v22)) + expected_dist = np.linalg.norm([H.dist(v11, v21), H.dist(v12, v22)], + ord='inf') assert dist == pytest.approx(expected_dist) diff --git a/odl/util/testutils.py b/odl/util/testutils.py index a86204fe139..c0ae330ed4b 100644 --- a/odl/util/testutils.py +++ b/odl/util/testutils.py @@ -207,7 +207,16 @@ def is_subdict(subdict, dictionary): def xfail_if(condition, reason=''): - """Return a ``pytest.xfail`` object if ``condition`` is ``True``.""" + """Return a ``pytest.xfail`` object if ``condition`` is ``True``. + + Examples + -------- + Create test that is expected to fail if ``condition`` is false, e.g. + + >>> condition = False + >>> with xfail_if(condition, reason='only works without condition'): + ... assert not condition + """ if condition: return pytest.xfail(reason) else: From 958dacc65f97e0c84295b61614748ee19df29874 Mon Sep 17 00:00:00 2001 From: Jonas Adler Date: Thu, 28 Jun 2018 14:17:04 +0200 Subject: [PATCH 38/38] WIP: Fixes to tests due to cupy rebase --- odl/contrib/mrc/test/uncompr_bin_test.py | 4 +- odl/test/space/tensors_test.py | 57 ++++++++++++------------ odl/util/utility.py | 15 ++----- 3 files changed, 34 insertions(+), 42 deletions(-) diff --git a/odl/contrib/mrc/test/uncompr_bin_test.py b/odl/contrib/mrc/test/uncompr_bin_test.py index 941f5189a9d..8219e8f2318 100644 --- a/odl/contrib/mrc/test/uncompr_bin_test.py +++ b/odl/contrib/mrc/test/uncompr_bin_test.py @@ -35,9 +35,9 @@ # --- Tests --- # -def test_uncompr_bin_io_without_header(shape, floating_dtype, order): +def test_uncompr_bin_io_without_header(shape, odl_floating_dtype, order): """Test I/O bypassing the header processing.""" - dtype = np.dtype(floating_dtype) + dtype = np.dtype(odl_floating_dtype) with tempfile.NamedTemporaryFile() as named_file: file = named_file.file diff --git a/odl/test/space/tensors_test.py b/odl/test/space/tensors_test.py index 8c8a1a6b1bb..2fde9b84b19 100644 --- a/odl/test/space/tensors_test.py +++ b/odl/test/space/tensors_test.py @@ -118,13 +118,14 @@ def weight(request): @pytest.fixture(scope='module') -def tspace(floating_dtype, odl_tspace_impl): +def tspace(odl_floating_dtype, odl_tspace_impl): available_dtypes = tensor_space_impl(odl_tspace_impl).available_dtypes() - if floating_dtype not in available_dtypes: + if odl_floating_dtype not in available_dtypes: pytest.skip('dtype {} not supported by impl {!r}' - ''.format(floating_dtype, odl_tspace_impl)) + ''.format(odl_floating_dtype, odl_tspace_impl)) else: - return odl.tensor_space(shape=(3, 4), dtype=floating_dtype, + return odl.tensor_space(shape=(3, 4), + dtype=odl_floating_dtype, impl=odl_tspace_impl) @@ -867,7 +868,7 @@ def test_element_setitem(odl_tspace_impl, setitem_indices): assert all_equal(x, x_arr) # Setting values with arrays - rhs_arr = array_module(tspace_impl).ones(sliced_shape) + rhs_arr = array_module(impl).ones(sliced_shape) x_arr[setitem_indices] = rhs_arr x[setitem_indices] = rhs_arr assert all_equal(x, x_arr) @@ -932,11 +933,12 @@ def test_element_setitem_bool_array(odl_tspace_impl): @skip_if_no_cupy -def test_asarray_numpy_to_cupy(floating_dtype): +def test_asarray_numpy_to_cupy(odl_floating_dtype): """Test x.asarray with numpy x and cupy impl and out.""" - space = odl.tensor_space((2, 3), dtype=floating_dtype) + dtype = odl_floating_dtype + space = odl.tensor_space((2, 3), dtype=dtype) - with xfail_if(floating_dtype in ('float128', 'complex256'), + with xfail_if(dtype in ('float128', 'complex256'), reason='quad precision types not available in cupy'): # Make new array, contiguous x = space.one() @@ -945,13 +947,13 @@ def test_asarray_numpy_to_cupy(floating_dtype): assert all_equal(x_cpy, x) # Write to existing, contiguous - out_cpy = cupy.empty((2, 3), dtype=floating_dtype) + out_cpy = cupy.empty((2, 3), dtype=dtype) x_cpy = x.asarray(out=out_cpy) assert x_cpy is out_cpy assert all_equal(out_cpy, x) # Make new array, discontiguous - arr = np.arange(12).astype(floating_dtype).reshape((2, 6))[:, ::2] + arr = np.arange(12).astype(dtype).reshape((2, 6))[:, ::2] x = space.element(arr) assert not (x.data.flags.c_contiguous or x.data.flags.f_contiguous) x_cpy = x.asarray(impl='cupy') @@ -959,18 +961,19 @@ def test_asarray_numpy_to_cupy(floating_dtype): assert all_equal(x_cpy, x) # Write to existing, contiguous - out_cpy = cupy.empty((2, 3), dtype=floating_dtype) + out_cpy = cupy.empty((2, 3), dtype=dtype) x_cpy = x.asarray(out=out_cpy) assert x_cpy is out_cpy assert all_equal(out_cpy, x) @skip_if_no_cupy -def test_asarray_cupy_to_numpy(floating_dtype): +def test_asarray_cupy_to_numpy(odl_floating_dtype): """Test x.asarray with cupy x and numpy impl and out.""" - with xfail_if(floating_dtype in ('float128', 'complex256'), + dtype = odl_floating_dtype + with xfail_if(dtype in ('float128', 'complex256'), reason='quad precision types not available in cupy'): - space = odl.tensor_space((2, 3), dtype=floating_dtype, impl='cupy') + space = odl.tensor_space((2, 3), dtype=dtype, impl='cupy') # Make new array, contiguous x = space.one() @@ -979,13 +982,13 @@ def test_asarray_cupy_to_numpy(floating_dtype): assert all_equal(x_npy, x) # Write to existing, contiguous - out_npy = np.empty((2, 3), dtype=floating_dtype) + out_npy = np.empty((2, 3), dtype=dtype) x_npy = x.asarray(out=out_npy) assert x_npy is out_npy assert all_equal(out_npy, x) # Make new array, discontiguous - arr = cupy.arange(12).astype(floating_dtype).reshape((2, 6))[:, ::2] + arr = cupy.arange(12).astype(dtype).reshape((2, 6))[:, ::2] x = space.element(arr) assert not (x.data.flags.c_contiguous or x.data.flags.f_contiguous) x_npy = x.asarray(impl='numpy') @@ -993,7 +996,7 @@ def test_asarray_cupy_to_numpy(floating_dtype): assert all_equal(x_npy, x) # Write to existing, contiguous - out_npy = np.empty((2, 3), dtype=floating_dtype) + out_npy = np.empty((2, 3), dtype=dtype) x_npy = x.asarray(out=out_npy) assert x_npy is out_npy assert all_equal(out_npy, x) @@ -1587,14 +1590,19 @@ def _check_result_is_out(result, out_seq): arrays, elements = noise_elements(tspace, nin + nout) # Arrays of the space's own data storage type in_arrays_own = arrays[:nin] - in_arrays_npy = [as_numpy(arr) for arr in arrays[:nin]] + in_arrays_npy = [as_numpy(arr) for arr in in_arrays_own] + out_arrays_own = arrays[:nin] + out_arrays_npy = [as_numpy(arr) for arr in out_arrays_own] data_elem = elements[0] + out_elems = elements[nin:] if nout == 1: - out_arr_kwargs = {'out': out_arrays[0]} + out_arr_own_kwargs = {'out': out_arrays_own[0]} + out_arr_npy_kwargs = {'out': out_arrays_npy[0]} out_elem_kwargs = {'out': out_elems[0]} elif nout > 1: - out_arr_kwargs = {'out': out_arrays[:nout]} + out_arr_own_kwargs = {'out': out_arrays_own[:nout]} + out_arr_npy_kwargs = {'out': out_arrays_npy[:nout]} out_elem_kwargs = {'out': out_elems[:nout]} # Get function to call, using both interfaces: @@ -1643,14 +1651,7 @@ def _check_result_is_out(result, out_seq): if USE_ARRAY_UFUNCS_INTERFACE: # Custom objects not allowed as `out` for numpy < 1.13 - result = ufunc_npy(*in_elems_npy, **out_elem_kwargs) - - if nout == 1: - kwargs_out = {'out': out_elems[0]} - elif nout == 2: - kwargs_out = {'out': (out_elems[0], out_elems[1])} - - result = ufunc_npy(*in_elems_npy, **kwargs_out) + result = ufunc_npy(*in_elems_npy, **out_arr_npy_kwargs) assert all_almost_equal(result_npy, result) _check_result_is_out(result, out_elems[:nout]) diff --git a/odl/util/utility.py b/odl/util/utility.py index 73e2fbe5137..e5a0422de69 100644 --- a/odl/util/utility.py +++ b/odl/util/utility.py @@ -21,18 +21,6 @@ import numpy as np __all__ = ( - 'array_str', 'dtype_str', 'dtype_repr', 'npy_printoptions', - 'signature_string', 'signature_string_parts', 'repr_string', - 'indent', 'dedent', 'attribute_repr_string', 'method_repr_string', - 'is_numeric_dtype', 'is_int_dtype', 'is_floating_dtype', 'is_real_dtype', - 'is_real_floating_dtype', 'is_complex_floating_dtype', - 'real_dtype', 'complex_dtype', 'is_string', 'nd_iterator', 'conj_exponent', - 'writable_array', 'run_from_ipython', 'NumpyRandomSeed', - 'cache_arguments', 'unique', - 'REPR_PRECISION') - -__all__ = ( - 'NumpyRandomSeed', 'array_cls', 'array_module', 'array_str', @@ -51,9 +39,12 @@ 'is_real_dtype', 'is_real_floating_dtype', 'is_string', + 'nd_iterator', 'none_context', 'npy_printoptions', + 'NumpyRandomSeed', 'real_dtype', + 'REPR_PRECISION', 'run_from_ipython', 'signature_string', 'unique',