diff --git a/.github/workflows/array-api-skips.txt b/.github/workflows/array-api-skips.txt index 2c94f9180fd4..fb81a683e6d7 100644 --- a/.github/workflows/array-api-skips.txt +++ b/.github/workflows/array-api-skips.txt @@ -1,7 +1,4 @@ # array API tests to be skipped -# data-apis/array-api-tests/issues/456 -array_api_tests/test_dlpack.py::test_dunder_dlpack - # data-apis/array-api/issues/1006 array_api_tests/test_special_cases.py::test_unary[tanh(real(x_i) is +infinity and isfinite(imag(x_i)) and imag(x_i) > 0) -> 1 + 0j] diff --git a/CHANGELOG.md b/CHANGELOG.md index 09961c7ccd34..1b1f37d76298 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -47,6 +47,8 @@ This release is compatible with NumPy 2.5. * Linked the `dpnp_backend_c` library against only the MKL SYCL domains it uses (`BLAS`, `RNG`, `VM`) [#3012](https://github.com/IntelPython/dpnp/pull/3012) * `dpnp` uses pybind11 3.1.0 [#3015](https://github.com/IntelPython/dpnp/pull/3015) * Reworked the ASV benchmarks and added end-to-end workload benchmarks derived from dpBench [#2996](https://github.com/IntelPython/dpnp/pull/2996) +* Reduced allocations in `dpnp.linalg.norm` by reusing the reduction result as the `sqrt` output buffer in the 2-norm and Frobenius-norm branches [#3062](https://github.com/IntelPython/dpnp/pull/3062) +* Avoided a copy of `dpnp.einsum` result into C-order by building the product in the requested layout directly [#3069](https://github.com/IntelPython/dpnp/pull/3069) ### Deprecated @@ -95,6 +97,9 @@ This release is compatible with NumPy 2.5. * Fixed `dpnp.insert` silently ignoring out-of-bounds negative indices in a multi-element `obj`, so a mix of in-bounds and out-of-bounds indices now consistently raises `IndexError` [#3041](https://github.com/IntelPython/dpnp/pull/3041) * Fixed a per-call `sycl::queue` leak in `usm_ndarray::get_queue()`/`get_device()` [#3042](https://github.com/IntelPython/dpnp/pull/3042) * Fixed `dpnp.linspace` returning `nan` for equal infinite endpoints [#3043](https://github.com/IntelPython/dpnp/pull/3043) +* Fixed `dpnp.cumsum`, `dpnp.cumprod`, and their `nan`/`cumulative_*` variants (including `dpnp.tensor.cumulative_sum`/`cumulative_prod`) silently returning incorrect results when accumulating along an axis of an array with more than one row [#3063](https://github.com/IntelPython/dpnp/pull/3063) +* Fixed `dpnp.einsum` returning a result whose memory layout differs from NumPy for the default `order="K"`, and ignoring `out` and `order` for a contraction over a size-0 dimension [#3058](https://github.com/IntelPython/dpnp/pull/3058) +* Fixed operations on a boolean array whose bytes are not `0x00`/`0x01` [#3055](https://github.com/IntelPython/dpnp/pull/3055) ### Security diff --git a/doc/comparison_generator.py b/doc/comparison_generator.py index 35190bec2f27..acb0dea2e9c6 100644 --- a/doc/comparison_generator.py +++ b/doc/comparison_generator.py @@ -38,21 +38,21 @@ def import_mod(mod, cls): obj = importlib.import_module(mod) if cls: obj = getattr(obj, cls) - return obj, ":meth:`{}.{}.{{}}`".format(mod, cls) + return obj, f":meth:`{mod}.{cls}.{{}}`" else: # ufunc is not a function - return obj, ":obj:`{}.{{}}`".format(mod) + return obj, f":obj:`{mod}.{{}}`" def generate_totals(base_mod, ref_mods, base_type, ref_types, cls): all_types = [base_type] + ref_types - header = ", ".join("**{} Total**".format(t) for t in all_types) - header = " {}".format(header) + header = ", ".join(f"**{t} Total**" for t in all_types) + header = f" {header}" totals = calc_totals(base_mod, ref_mods, cls) cells = ", ".join(str(t) for t in totals) - total = " {}".format(cells) + total = f" {cells}" return [header, total] @@ -79,14 +79,12 @@ def generate_comparison_rst(base_mod, ref_mods, base_type, ref_types, cls): ref_cells.append(ref_cell) cells = ", ".join([base_cell] + ref_cells) - line = " {}".format(cells) + line = f" {cells}" rows.append(line) totals = generate_totals(base_mod, ref_mods, base_type, ref_types, cls) - return ( - [".. csv-table::", " :header: {}".format(header), ""] + rows + totals - ) + return [".. csv-table::", f" :header: {header}", ""] + rows + totals def section(header, base_mod, ref_mods, base_type, ref_types, cls=None): @@ -111,15 +109,15 @@ def generate_totals_numbers(header, base_mod, ref_mods, cls=None): totals = [header] + calc_totals(base_mod, ref_mods, cls) cells = ", ".join(str(t) for t in totals) - total = " {}".format(cells) + total = f" {cells}" return total, counter_funcs def generate_table_numbers(base_mod, ref_mods, base_type, ref_types, cls=None): all_types = ["Name"] + [base_type] + ref_types - header = ", ".join("**{}**".format(t) for t in all_types) - header = " {}".format(header) + header = ", ".join(f"**{t}**" for t in all_types) + header = f" {header}" rows = [] counters_funcs = [] @@ -131,7 +129,7 @@ def generate_table_numbers(base_mod, ref_mods, base_type, ref_types, cls=None): totals.append(totals_) counters_funcs.append(counters_funcs_) cells = ", ".join(str(t) for t in totals) - total = " {}".format(cells) + total = f" {cells}" rows.append(total) totals = [] @@ -141,7 +139,7 @@ def generate_table_numbers(base_mod, ref_mods, base_type, ref_types, cls=None): totals.append(totals_) counters_funcs.append(counters_funcs_) cells = ", ".join(str(t) for t in totals) - total = " {}".format(cells) + total = f" {cells}" rows.append(total) totals = [] @@ -153,7 +151,7 @@ def generate_table_numbers(base_mod, ref_mods, base_type, ref_types, cls=None): totals.append(totals_) counters_funcs.append(counters_funcs_) cells = ", ".join(str(t) for t in totals) - total = " {}".format(cells) + total = f" {cells}" rows.append(total) totals = [] @@ -165,7 +163,7 @@ def generate_table_numbers(base_mod, ref_mods, base_type, ref_types, cls=None): totals.append(totals_) counters_funcs.append(counters_funcs_) cells = ", ".join(str(t) for t in totals) - total = " {}".format(cells) + total = f" {cells}" rows.append(total) totals = [] @@ -177,7 +175,7 @@ def generate_table_numbers(base_mod, ref_mods, base_type, ref_types, cls=None): totals.append(totals_) counters_funcs.append(counters_funcs_) cells = ", ".join(str(t) for t in totals) - total = " {}".format(cells) + total = f" {cells}" rows.append(total) counter_functions = [] @@ -185,11 +183,11 @@ def generate_table_numbers(base_mod, ref_mods, base_type, ref_types, cls=None): counter = 0 for j in range(len(counters_funcs)): counter += counters_funcs[j][i] - counter_functions.append("{}".format(counter)) + counter_functions.append(f"{counter}") summary = ["Total"] + counter_functions cells = ", ".join(str(t) for t in summary) - summary_total = " {}".format(cells) + summary_total = f" {cells}" rows.append(summary_total) comparison_rst = [".. csv-table::", ""] + [header] + rows @@ -207,7 +205,7 @@ def generate(): ref_mods += ["dpnp"] ref_types += ["DPNP"] - ref_vers = ["DPNP(v{})".format(dpnp.__version__)] + ref_vers = [f"DPNP(v{dpnp.__version__})"] except ImportError as err: print(f"DOCBUILD: Can't load DPNP module with error={err}") @@ -216,7 +214,7 @@ def generate(): ref_mods += ["cupy"] ref_types += ["CuPy"] - ref_vers += ["CuPy(v{})".format(cupy.__version__)] + ref_vers += [f"CuPy(v{cupy.__version__})"] except ImportError as err: print(f"DOCBUILD: Can't load CuPy module with error={err}") @@ -225,12 +223,12 @@ def generate(): base_mod = "numpy" # TODO: Why string? base_type = "NumPy" - base_ver = "{}(v{})".format(base_type, numpy.__version__) + base_ver = f"{base_type}(v{numpy.__version__})" except ImportError as err: print(f"DOCBUILD: Can't load {base_type} module with error={err}") header = " / ".join([base_ver] + ref_vers) + " APIs" - buf = ["**{}**".format(header), ""] + buf = [f"**{header}**", ""] buf += generate_table_numbers(base_mod, ref_mods, base_type, ref_types) buf += section("Module-Level", base_mod, ref_mods, base_type, ref_types) diff --git a/dpnp/dpnp_array.py b/dpnp/dpnp_array.py index 86055a4828fd..b225fb2c7329 100644 --- a/dpnp/dpnp_array.py +++ b/dpnp/dpnp_array.py @@ -160,15 +160,15 @@ def __init__( ) def __abs__(self, /): - r"""Return :math:`|\text{self}|`.""" + """Return ``|self|``.""" return dpnp.abs(self) def __add__(self, other, /): - r"""Return :math:`\text{self + value}`.""" + """Return ``self + value``.""" return dpnp.add(self, other) def __and__(self, other, /): - r"""Return :math:`\text{self & value}`.""" + """Return ``self & value``.""" return dpnp.bitwise_and(self, other) def __array__(self, dtype=None, /, *, copy=None): @@ -228,7 +228,7 @@ def __bool__(self, /): return self._array_obj.__bool__() def __bytes__(self): - r"""Return :math:`\text{bytes(self)}`.""" + """Return ``bytes(self)``.""" return bytes(self.asnumpy()) # '__class__', @@ -239,7 +239,7 @@ def __complex__(self, /): return self._array_obj.__complex__() def __contains__(self, value, /): - r"""Return :math:`\text{value in self}`.""" + """Return ``value in self``.""" return (self == value).any() def __copy__(self): @@ -256,7 +256,7 @@ def __copy__(self): # '__dir__', def __divmod__(self, other, /): - r"""Return :math:`\text{divmod(self, value)}`.""" + """Return ``divmod(self, value)``.""" return dpnp.divmod(self, other) def __dlpack__( @@ -340,7 +340,7 @@ def __dlpack_device__(self, /): return self._array_obj.__dlpack_device__() def __eq__(self, other, /): - r"""Return :math:`\text{self == value}`.""" + """Return ``self == value``.""" return dpnp.equal(self, other) def __float__(self, /): @@ -348,19 +348,19 @@ def __float__(self, /): return self._array_obj.__float__() def __floordiv__(self, other, /): - r"""Return :math:`\text{self // value}`.""" + """Return ``self // value``.""" return dpnp.floor_divide(self, other) def __format__(self, format_spec): - r"""Return :math:`\text{format(self, format_spec)}`.""" + """Return ``format(self, format_spec)``.""" return format(self.asnumpy(), format_spec) def __ge__(self, other, /): - r"""Return :math:`\text{self >= value}`.""" + """Return ``self >= value``.""" return dpnp.greater_equal(self, other) def __getitem__(self, key, /): - r"""Return :math:`\text{self[key]}`.""" + """Return ``self[key]``.""" key = _get_unwrapped_index_key(key) item = self._array_obj.__getitem__(key) @@ -369,33 +369,33 @@ def __getitem__(self, key, /): # '__getstate__', def __gt__(self, other, /): - r"""Return :math:`\text{self > value}`.""" + """Return ``self > value``.""" return dpnp.greater(self, other) # '__hash__', def __iadd__(self, other, /): - r"""Return :math:`\text{self += value}`.""" + """Return ``self += value``.""" dpnp.add(self, other, out=self) return self def __iand__(self, other, /): - r"""Return :math:`\text{self &= value}`.""" + """Return ``self &= value``.""" dpnp.bitwise_and(self, other, out=self) return self def __ifloordiv__(self, other, /): - r"""Return :math:`\text{self //= value}`.""" + """Return ``self //= value``.""" dpnp.floor_divide(self, other, out=self) return self def __ilshift__(self, other, /): - r"""Return :math:`\text{self <<= value}`.""" + """Return ``self <<= value``.""" dpnp.left_shift(self, other, out=self) return self def __imatmul__(self, other, /): - r"""Return :math:`\text{self @= value}`.""" + """Return ``self @= value``.""" # Unlike `matmul(a, b, out=a)` we ensure that the result isn't broadcast # if the result without `out` would have less dimensions than `a`. @@ -419,12 +419,12 @@ def __imatmul__(self, other, /): return self def __imod__(self, other, /): - r"""Return :math:`\text{self %= value}`.""" + """Return ``self %= value``.""" dpnp.remainder(self, other, out=self) return self def __imul__(self, other, /): - r"""Return :math:`\text{self *= value}`.""" + """Return ``self *= value``.""" dpnp.multiply(self, other, out=self) return self @@ -439,168 +439,168 @@ def __int__(self, /): return self._array_obj.__int__() def __invert__(self, /): - r"""Return :math:`\text{~self}`.""" + """Return ``~self``.""" return dpnp.invert(self) def __ior__(self, other, /): - r"""Return :math:`\text{self |= value}`.""" + """Return ``self |= value``.""" dpnp.bitwise_or(self, other, out=self) return self def __ipow__(self, other, /): - r"""Return :math:`\text{self **= value}`.""" + """Return ``self **= value``.""" dpnp.power(self, other, out=self) return self def __irshift__(self, other, /): - r"""Return :math:`\text{self >>= value}`.""" + """Return ``self >>= value``.""" dpnp.right_shift(self, other, out=self) return self def __isub__(self, other, /): - r"""Return :math:`\text{self -= value}`.""" + """Return ``self -= value``.""" dpnp.subtract(self, other, out=self) return self def __iter__(self, /): - r"""Return :math:`\text{iter(self)}`.""" + """Return ``iter(self)``.""" if self.ndim == 0: raise TypeError("iteration over a 0-d array") return (self[i] for i in range(self.shape[0])) def __itruediv__(self, other, /): - r"""Return :math:`\text{self /= value}`.""" + """Return ``self /= value``.""" dpnp.true_divide(self, other, out=self) return self def __ixor__(self, other, /): - r"""Return :math:`\text{self ^= value}`.""" + """Return ``self ^= value``.""" dpnp.bitwise_xor(self, other, out=self) return self def __le__(self, other, /): - r"""Return :math:`\text{self <= value}`.""" + """Return ``self <= value``.""" return dpnp.less_equal(self, other) def __len__(self): - r"""Return :math:`\text{len(self)}`.""" + """Return ``len(self)``.""" return self._array_obj.__len__() def __lshift__(self, other, /): - r"""Return :math:`\text{self << value}`.""" + """Return ``self << value``.""" return dpnp.left_shift(self, other) def __lt__(self, other, /): - r"""Return :math:`\text{self < value}`.""" + """Return ``self < value``.""" return dpnp.less(self, other) def __matmul__(self, other, /): - r"""Return :math:`\text{self @ value}`.""" + """Return ``self @ value``.""" return dpnp.matmul(self, other) def __mod__(self, other, /): - r"""Return :math:`\text{self % value}`.""" + """Return ``self % value``.""" return dpnp.remainder(self, other) def __mul__(self, other, /): - r"""Return :math:`\text{self * value}`.""" + """Return ``self * value``.""" return dpnp.multiply(self, other) def __ne__(self, other, /): - r"""Return :math:`\text{self != value}`.""" + """Return ``self != value``.""" return dpnp.not_equal(self, other) def __neg__(self, /): - r"""Return :math:`\text{-self}`.""" + """Return ``-self``.""" return dpnp.negative(self) # '__new__', def __or__(self, other, /): - r"""Return :math:`\text{self | value}`.""" + """Return ``self | value``.""" return dpnp.bitwise_or(self, other) def __pos__(self, /): - r"""Return :math:`\text{+self}`.""" + """Return ``+self``.""" return dpnp.positive(self) def __pow__(self, other, mod=None, /): - r"""Return :math:`\text{self ** value}`.""" + """Return ``self ** value``.""" if mod is not None: return NotImplemented return dpnp.power(self, other) def __radd__(self, other, /): - r"""Return :math:`\text{value + self}`.""" + """Return ``value + self``.""" return dpnp.add(other, self) def __rand__(self, other, /): - r"""Return :math:`\text{value & self}`.""" + """Return ``value & self``.""" return dpnp.bitwise_and(other, self) def __rdivmod__(self, other, /): - r"""Return :math:`\text{divmod(value, self)}`.""" + """Return ``divmod(value, self)``.""" return dpnp.divmod(other, self) # '__reduce__', # '__reduce_ex__', def __repr__(self): - r"""Return :math:`\text{repr(self)}`.""" + """Return ``repr(self)``.""" return dpt.usm_ndarray_repr(self._array_obj, prefix="array") def __rfloordiv__(self, other, /): - r"""Return :math:`\text{value // self}`.""" + """Return ``value // self``.""" return dpnp.floor_divide(other, self) def __rlshift__(self, other, /): - r"""Return :math:`\text{value << self}`.""" + """Return ``value << self``.""" return dpnp.left_shift(other, self) def __rmatmul__(self, other, /): - r"""Return :math:`\text{value @ self}`.""" + """Return ``value @ self``.""" return dpnp.matmul(other, self) def __rmod__(self, other, /): - r"""Return :math:`\text{value % self}`.""" + """Return ``value % self``.""" return dpnp.remainder(other, self) def __rmul__(self, other, /): - r"""Return :math:`\text{value * self}`.""" + """Return ``value * self``.""" return dpnp.multiply(other, self) def __ror__(self, other, /): - r"""Return :math:`\text{value | self}`.""" + """Return ``value | self``.""" return dpnp.bitwise_or(other, self) def __rpow__(self, other, mod=None, /): - r"""Return :math:`\text{value ** self}`.""" + """Return ``value ** self``.""" if mod is not None: return NotImplemented return dpnp.power(other, self) def __rrshift__(self, other, /): - r"""Return :math:`\text{value >> self}`.""" + """Return ``value >> self``.""" return dpnp.right_shift(other, self) def __rshift__(self, other, /): - r"""Return :math:`\text{self >> value}`.""" + """Return ``self >> value``.""" return dpnp.right_shift(self, other) def __rsub__(self, other, /): - r"""Return :math:`\text{value - self}`.""" + """Return ``value - self``.""" return dpnp.subtract(other, self) def __rtruediv__(self, other, /): - r"""Return :math:`\text{value / self}`.""" + """Return ``value / self``.""" return dpnp.true_divide(other, self) def __rxor__(self, other, /): - r"""Return :math:`\text{value ^ self}`.""" + """Return ``value ^ self``.""" return dpnp.bitwise_xor(other, self) def __setitem__(self, key, value, /): - r"""Set :math:`\text{self[key]}` to a value.""" + """Set ``self[key]`` to a value.""" key = _get_unwrapped_index_key(key) if isinstance(value, dpnp_array): @@ -614,11 +614,11 @@ def __setitem__(self, key, value, /): __slots__ = ("_array_obj",) def __str__(self): - r"""Return :math:`\text{str(self)}`.""" + """Return ``str(self)``.""" return self._array_obj.__str__() def __sub__(self, other, /): - r"""Return :math:`\text{self - value}`.""" + """Return ``self - value``.""" return dpnp.subtract(self, other) @property @@ -630,7 +630,7 @@ def __sycl_usm_array_interface__(self): return self._array_obj.__sycl_usm_array_interface__ def __truediv__(self, other, /): - r"""Return :math:`\text{self / value}`.""" + """Return ``self / value``.""" return dpnp.true_divide(self, other) @property @@ -653,7 +653,7 @@ def __usm_ndarray__(self): return self._array_obj def __xor__(self, other, /): - r"""Return :math:`\text{self ^ value}`.""" + """Return ``self ^ value``.""" return dpnp.bitwise_xor(self, other) @staticmethod diff --git a/dpnp/dpnp_flatiter.py b/dpnp/dpnp_flatiter.py index 7375e03d8020..bc17cdd22ae2 100644 --- a/dpnp/dpnp_flatiter.py +++ b/dpnp/dpnp_flatiter.py @@ -37,7 +37,7 @@ class flatiter: def __init__(self, X): if type(X) is not dpnp.ndarray: raise TypeError( - "Argument must be of type dpnp.ndarray, got {}".format(type(X)) + f"Argument must be of type dpnp.ndarray, got {type(X)}" ) self.arr_ = X self.size_ = X.size diff --git a/dpnp/dpnp_iface_bitwise.py b/dpnp/dpnp_iface_bitwise.py index 604fd365ee18..350cc740782d 100644 --- a/dpnp/dpnp_iface_bitwise.py +++ b/dpnp/dpnp_iface_bitwise.py @@ -550,6 +550,7 @@ def binary_repr(num, width=None): Memory layout of the newly output array, if parameter `out` is ``None``. Default: ``"K"``. + Returns ------- out : dpnp.ndarray diff --git a/dpnp/dpnp_iface_histograms.py b/dpnp/dpnp_iface_histograms.py index 8cd8bd3c08e8..1ac51d18f4c8 100644 --- a/dpnp/dpnp_iface_histograms.py +++ b/dpnp/dpnp_iface_histograms.py @@ -574,7 +574,7 @@ def histogram(a, bins=10, range=None, density=None, weights=None): :obj:`dpnp.digitize` : Return the indices of the bins to which each value in input array belongs. :obj:`dpnp.histogram_bin_edges` : Return only the edges of the bins used - by the obj:`dpnp.histogram` function. + by the :obj:`dpnp.histogram` function. Examples -------- @@ -825,6 +825,7 @@ def histogram2d(x, y, bins=10, range=None, density=None, weights=None): If ``None`` all samples are assigned a weight of ``1``. Default: ``None``. + Returns ------- H : dpnp.ndarray of shape (nx, ny) diff --git a/dpnp/dpnp_iface_sorting.py b/dpnp/dpnp_iface_sorting.py index 8f6f3e80f0d1..5feffffbc432 100644 --- a/dpnp/dpnp_iface_sorting.py +++ b/dpnp/dpnp_iface_sorting.py @@ -156,7 +156,7 @@ def argsort( :obj:`dpnp.sort` : Return a sorted copy of an array. :obj:`dpnp.lexsort` : Indirect stable sort with multiple keys. :obj:`dpnp.argpartition` : Indirect partial sort. - :obj:`dpnp.take_along_axis` : Apply ``index_array`` from obj:`dpnp.argsort` + :obj:`dpnp.take_along_axis` : Apply ``index_array`` from :obj:`dpnp.argsort` to an array as if by calling sort. Examples diff --git a/dpnp/dpnp_utils/dpnp_algo_utils.pyx b/dpnp/dpnp_utils/dpnp_algo_utils.pyx index d13de1c5cdb9..5a3fdd397501 100644 --- a/dpnp/dpnp_utils/dpnp_algo_utils.pyx +++ b/dpnp/dpnp_utils/dpnp_algo_utils.pyx @@ -304,9 +304,7 @@ def get_usm_allocations(objects): if not isinstance(objects, (list, tuple)): raise TypeError( - "Expected a list or a tuple, got {}".format( - type(objects) - ) + f"Expected a list or a tuple, got {type(objects)}" ) if len(objects) == 0: @@ -517,11 +515,8 @@ cdef tuple get_common_usm_allocation( if common_usm_type is None: raise ValueError( "could not recognize common USM type " - "for inputs of USM types {} and {}" - "".format( - array1_obj.usm_type, - array2_obj.usm_type, - ) + f"for inputs of USM types {array1_obj.usm_type} " + f"and {array2_obj.usm_type}" ) common_sycl_queue = get_execution_queue( @@ -530,11 +525,8 @@ cdef tuple get_common_usm_allocation( if common_sycl_queue is None: raise ValueError( "could not recognize common SYCL queue " - "for inputs in SYCL queues {} and {}" - "".format( - array1_obj.sycl_queue, - array2_obj.sycl_queue, - ) + f"for inputs in SYCL queues {array1_obj.sycl_queue} " + f"and {array2_obj.sycl_queue}" ) return ( @@ -682,8 +674,7 @@ cdef class dpnp_descriptor: raise TypeError( "expected either dpnp.tensor.usm_ndarray " - "or dpnp.dpnp_array.dpnp_array, got {}" - "".format(type(self.origin_pyobj)) + f"or dpnp.dpnp_array.dpnp_array, got {type(self.origin_pyobj)}" ) cdef void * get_data(self): diff --git a/dpnp/dpnp_utils/dpnp_utils_einsum.py b/dpnp/dpnp_utils/dpnp_utils_einsum.py index 842596bdc32d..1f2521ef8232 100644 --- a/dpnp/dpnp_utils/dpnp_utils_einsum.py +++ b/dpnp/dpnp_utils/dpnp_utils_einsum.py @@ -880,7 +880,7 @@ def _parse_possible_contraction( return [sort, positions, new_input_sets] -def _reduced_binary_einsum(arr0, sub0, arr1, sub1, sub_others): +def _reduced_binary_einsum(arr0, sub0, arr1, sub1, sub_others, prefer_c=False): """Copied from _reduced_binary_einsum in cupy/core/_einsum.py""" set0 = set(sub0) @@ -916,6 +916,27 @@ def _reduced_binary_einsum(arr0, sub0, arr1, sub1, sub_others): arr1 = _expand_dims_transpose(arr1, sub1, sub_out) return arr0 * arr1, sub_out + if ( + prefer_c + and sub_l + and sub_r + and sub_others.index(sub_l[0]) > sub_others.index(sub_r[0]) + ): + # Swap the operand roles so the free axes of the product already come + # out in `sub_others` order. Otherwise the caller transposes the + # result, which makes it f-contiguous and costs a copy into c-order. + sub_out = sub_b + sub_r + sub_l + arr0, bs0, cs0, ts0, arr1, bs1, cs1, ts1 = ( + arr1, + bs1, + cs1, + ts1, + arr0, + bs0, + cs0, + ts0, + ) + tmp0, shapes0 = _flatten_transpose(arr0, [bs0, ts0, cs0]) tmp1, shapes1 = _flatten_transpose(arr1, [bs1, cs1, ts1]) shapes_out = shapes0[0] + shapes0[1] + shapes1[2] @@ -1039,8 +1060,22 @@ def dpnp_einsum( ) arrays.append(operands[id]) result_dtype = dpnp.result_type(*arrays) if dtype is None else dtype - if order is not None and order in "aA": - order = "F" if all(arr.flags.fnc for arr in arrays) else "C" + # validated here because the view path below skips `dpnp.asarray` + if order is None: + order = "K" + elif not isinstance(order, str): + raise TypeError(f"order must be str, not {type(order).__name__}") + elif len(order) == 1 and order in "afkcAFKC": + order = order.upper() + else: + raise ValueError( + f"order must be one of 'C', 'F', 'A', or 'K' (got '{order}')" + ) + all_f_contiguous = all(arr.flags.f_contiguous for arr in arrays) + if order == "A": + # NumPy uses f_contiguous here, not fnc; they differ for an array that + # is both C- and F-contiguous, such as a 1-D or size-1 one + order = "F" if all_f_contiguous else "C" input_subscripts = [ _parse_ellipsis_subscript(sub, idx, ndim=arr.ndim) @@ -1110,12 +1145,15 @@ def dpnp_einsum( # no more raises if len(operands) >= 2: if any(arr.size == 0 for arr in operands): - return dpnp.zeros( + # NumPy falls back to "C" for "K" here + arr_out = dpnp.zeros( tuple(dimension_dict[label] for label in output_subscript), dtype=result_dtype, + order="C" if order == "K" else order, usm_type=res_usm_type, sycl_queue=exec_q, ) + return dpnp.get_result_array(arr_out, out, casting=casting) # Don't squeeze if unary, because this affects later (in trivial sum) # whether the return is a writeable view. @@ -1194,7 +1232,15 @@ def dpnp_einsum( stacklevel=2, ) - for idx0, idx1 in _iter_path_pairs(path): + # Resolved above the loop because `prefer_c` below needs the final order: + # only a "C" target gains from a product laid out in output order. + if order == "K" and optimize is False and not all_f_contiguous: + # only the unoptimized path of NumPy copies into a c-contiguous + # array, the optimized one is matmul-based, as dpnp always is + order = "C" + + pairs = list(_iter_path_pairs(path)) + for pair_idx, (idx0, idx1) in enumerate(pairs): # "reduced" binary einsum arr0 = operands.pop(idx0) sub0 = input_subscripts.pop(idx0) @@ -1207,7 +1253,13 @@ def dpnp_einsum( ) ) arr_out, sub_out = _reduced_binary_einsum( - arr0, sub0, arr1, sub1, sub_others + arr0, + sub0, + arr1, + sub1, + sub_others, + # only "C" and the last contraction + prefer_c=order == "C" and pair_idx == len(pairs) - 1, ) operands.append(arr_out) input_subscripts.append(sub_out) @@ -1226,6 +1278,8 @@ def dpnp_einsum( [dimension_dict[label] for label in output_subscript] ) - arr_out = dpnp.asarray(arr_out, order=order) + # a view is returned for any `order`, the same way NumPy does + if not returns_view: + arr_out = dpnp.asarray(arr_out, order=order) assert returns_view or arr_out.dtype == result_dtype return dpnp.get_result_array(arr_out, out, casting=casting) diff --git a/dpnp/linalg/dpnp_utils_linalg.py b/dpnp/linalg/dpnp_utils_linalg.py index 527235496b9d..ac1c87a61b5b 100644 --- a/dpnp/linalg/dpnp_utils_linalg.py +++ b/dpnp/linalg/dpnp_utils_linalg.py @@ -1162,8 +1162,8 @@ def _norm_int_axis(x, ord, axis, keepdims): return dpnp.abs(x).sum(axis=axis, keepdims=keepdims) if ord is None or ord == 2: # special case for speedup - s = (dpnp.conj(x) * x).real - return dpnp.sqrt(dpnp.sum(s, axis=axis, keepdims=keepdims)) + s = dpnp.sum((dpnp.conj(x) * x).real, axis=axis, keepdims=keepdims) + return dpnp.sqrt(s, out=s) if isinstance(ord, (int, float)): absx = dpnp.abs(x) absx **= ord @@ -1215,7 +1215,8 @@ def _norm_tuple_axis(x, ord, row_axis, col_axis, keepdims): row_axis -= 1 ret = dpnp.abs(x).sum(axis=col_axis).min(axis=row_axis) elif ord in [None, "fro", "f"]: - ret = dpnp.sqrt(dpnp.sum((dpnp.conj(x) * x).real, axis=axis)) + ret = dpnp.sum((dpnp.conj(x) * x).real, axis=axis) + ret = dpnp.sqrt(ret, out=ret) elif ord == "nuc": ret = _multi_svd_norm(x, row_axis, col_axis, dpnp.sum) else: @@ -2360,7 +2361,7 @@ def dpnp_norm(x, ord=None, axis=None, keepdims=False): sqnorm = dpnp.dot(x_real, x_real) + dpnp.dot(x_imag, x_imag) else: sqnorm = dpnp.dot(x, x) - ret = dpnp.sqrt(sqnorm) + ret = dpnp.sqrt(sqnorm, out=sqnorm) if keepdims: ret = ret.reshape((1,) * ndim) return ret diff --git a/dpnp/random/dpnp_iface_random.py b/dpnp/random/dpnp_iface_random.py index 5db780864280..ca1e3d1a463a 100644 --- a/dpnp/random/dpnp_iface_random.py +++ b/dpnp/random/dpnp_iface_random.py @@ -79,9 +79,8 @@ def _get_random_state(device=None, sycl_queue=None): _dpnp_random_states[sycl_queue] = rs else: raise RuntimeError( - "Normalized SYCL queue {} mismatched with one returned by RandmoState {}".format( - sycl_queue, rs.get_sycl_queue() - ) + f"Normalized SYCL queue {sycl_queue} mismatched with one " + f"returned by RandmoState {rs.get_sycl_queue()}" ) return _dpnp_random_states[sycl_queue] diff --git a/dpnp/random/dpnp_random_state.py b/dpnp/random/dpnp_random_state.py index a18148648d7f..d762c299c1f3 100644 --- a/dpnp/random/dpnp_random_state.py +++ b/dpnp/random/dpnp_random_state.py @@ -119,7 +119,7 @@ def __init__(self, seed=None, device=None, sycl_queue=None): ) def __repr__(self): - return self.__str__() + " at 0x{:X}".format(id(self)) + return self.__str__() + f" at 0x{id(self):X}" def __str__(self): _str = self.__class__.__name__ diff --git a/dpnp/tensor/_accumulation.py b/dpnp/tensor/_accumulation.py index 069eb870f783..4def71efc587 100644 --- a/dpnp/tensor/_accumulation.py +++ b/dpnp/tensor/_accumulation.py @@ -61,7 +61,7 @@ def _accumulate_common( if axis is None: if nd > 1: raise ValueError( - "`axis` cannot be `None` for array of dimension `{}`".format(nd) + f"`axis` cannot be `None` for array of dimension `{nd}`" ) axis = 0 else: diff --git a/dpnp/tensor/_compute_follows_data.pyx b/dpnp/tensor/_compute_follows_data.pyx index 798c158c7e1b..65523a48c2cb 100644 --- a/dpnp/tensor/_compute_follows_data.pyx +++ b/dpnp/tensor/_compute_follows_data.pyx @@ -80,7 +80,7 @@ def get_execution_queue(qs, /): """ if not isinstance(qs, (list, tuple)): raise TypeError( - "Expected a list or a tuple, got {}".format(type(qs)) + f"Expected a list or a tuple, got {type(qs)}" ) if len(qs) == 0: return None @@ -114,7 +114,7 @@ def get_coerced_usm_type(usm_types, /): """ if not isinstance(usm_types, (list, tuple)): raise TypeError( - "Expected a list or a tuple, got {}".format(type(usm_types)) + f"Expected a list or a tuple, got {type(usm_types)}" ) if len(usm_types) == 0: return None diff --git a/dpnp/tensor/_dlpack.pyx b/dpnp/tensor/_dlpack.pyx index 262970385adc..9e83f46de534 100644 --- a/dpnp/tensor/_dlpack.pyx +++ b/dpnp/tensor/_dlpack.pyx @@ -709,9 +709,7 @@ cdef dict _numpy_array_interface_from_dl_tensor(DLTensor *dlt, bint ro_flag): ary_dt = "b" + str(itemsize) else: raise BufferError( - "Can not import DLPack tensor with type code {}.".format( - dlt.dtype.code - ) + f"Can not import DLPack tensor with type code {dlt.dtype.code}." ) typestr = "|" + ary_dt return dict( @@ -937,9 +935,8 @@ cpdef object from_dlpack_capsule(object py_caps): ary_dt = np.dtype("?") else: raise BufferError( - "Can not import DLPack tensor with type code {}.".format( - dl_tensor.dtype.code - ) + "Can not import DLPack tensor with type code " + f"{dl_tensor.dtype.code}." ) res_ary = usm_ndarray( py_shape, diff --git a/dpnp/tensor/_flags.pyx b/dpnp/tensor/_flags.pyx index 07cae3f3e930..48adf80116cd 100644 --- a/dpnp/tensor/_flags.pyx +++ b/dpnp/tensor/_flags.pyx @@ -162,7 +162,7 @@ cdef class Flags: def __repr__(self): out = [] for name in "C_CONTIGUOUS", "F_CONTIGUOUS", "WRITABLE": - out.append(" {} : {}".format(name, self[name])) + out.append(f" {name} : {self[name]}") return "\n".join(out) def __eq__(self, other): diff --git a/dpnp/tensor/_indexing_functions.py b/dpnp/tensor/_indexing_functions.py index 48849ff7bbd1..b61edb892a0c 100644 --- a/dpnp/tensor/_indexing_functions.py +++ b/dpnp/tensor/_indexing_functions.py @@ -47,9 +47,7 @@ def _get_indexing_mode(name): try: return modes[name] except KeyError: - raise ValueError( - "`mode` must be `wrap` or `clip`." "Got `{}`.".format(name) - ) + raise ValueError(f"`mode` must be `wrap` or `clip`. Got `{name}`.") def _range(sh_i, i, nd, q, usm_t, dt): @@ -272,13 +270,11 @@ def put_vec_duplicates(vec, ind, vals): """ if not isinstance(x, dpt.usm_ndarray): raise TypeError( - "Expected instance of `dpt.usm_ndarray`, got `{}`.".format(type(x)) + f"Expected instance of `dpt.usm_ndarray`, got `{type(x)}`." ) if not isinstance(indices, dpt.usm_ndarray): raise TypeError( - "`indices` expected `dpt.usm_ndarray`, got `{}`.".format( - type(indices) - ) + f"`indices` expected `dpt.usm_ndarray`, got `{type(indices)}`." ) if isinstance(vals, dpt.usm_ndarray): queues_ = [x.sycl_queue, indices.sycl_queue, vals.sycl_queue] @@ -287,14 +283,10 @@ def put_vec_duplicates(vec, ind, vals): queues_ = [x.sycl_queue, indices.sycl_queue] usm_types_ = [x.usm_type, indices.usm_type] if indices.ndim != 1: - raise ValueError( - "`indices` expected a 1D array, got `{}`".format(indices.ndim) - ) + raise ValueError(f"`indices` expected a 1D array, got `{indices.ndim}`") if indices.dtype.kind not in "ui": raise IndexError( - "`indices` expected integer data type, got `{}`".format( - indices.dtype - ) + f"`indices` expected integer data type, got `{indices.dtype}`" ) exec_q = dpt.get_execution_queue(queues_) if exec_q is None: @@ -307,9 +299,7 @@ def put_vec_duplicates(vec, ind, vals): if axis is None: if x_ndim > 1: raise ValueError( - "`axis` cannot be `None` for array of dimension `{}`".format( - x_ndim - ) + f"`axis` cannot be `None` for array of dimension `{x_ndim}`" ) axis = 0 @@ -469,25 +459,19 @@ def take(x, indices, /, *, axis=None, out=None, mode="wrap"): """ if not isinstance(x, dpt.usm_ndarray): raise TypeError( - "Expected instance of `dpt.usm_ndarray`, got `{}`.".format(type(x)) + f"Expected instance of `dpt.usm_ndarray`, got `{type(x)}`." ) if not isinstance(indices, dpt.usm_ndarray): raise TypeError( - "`indices` expected `dpt.usm_ndarray`, got `{}`.".format( - type(indices) - ) + f"`indices` expected `dpt.usm_ndarray`, got `{type(indices)}`." ) if indices.dtype.kind not in "ui": raise IndexError( - "`indices` expected integer data type, got `{}`".format( - indices.dtype - ) + f"`indices` expected integer data type, got `{indices.dtype}`" ) if indices.ndim != 1: - raise ValueError( - "`indices` expected a 1D array, got `{}`".format(indices.ndim) - ) + raise ValueError(f"`indices` expected a 1D array, got `{indices.ndim}`") exec_q = dpt.get_execution_queue([x.sycl_queue, indices.sycl_queue]) if exec_q is None: raise dpt.ExecutionPlacementError @@ -499,9 +483,7 @@ def take(x, indices, /, *, axis=None, out=None, mode="wrap"): if axis is None: if x_ndim > 1: raise ValueError( - "`axis` cannot be `None` for array of dimension `{}`".format( - x_ndim - ) + f"`axis` cannot be `None` for array of dimension `{x_ndim}`" ) axis = 0 diff --git a/dpnp/tensor/_linear_algebra_functions.py b/dpnp/tensor/_linear_algebra_functions.py index dcaf99b4423c..c349784fdd83 100644 --- a/dpnp/tensor/_linear_algebra_functions.py +++ b/dpnp/tensor/_linear_algebra_functions.py @@ -65,7 +65,7 @@ def matrix_transpose(x): if not isinstance(x, dpt.usm_ndarray): raise TypeError( - "Expected instance of `dpt.usm_ndarray`, got `{}`.".format(type(x)) + f"Expected instance of `dpt.usm_ndarray`, got `{type(x)}`." ) if x.ndim < 2: raise ValueError( diff --git a/dpnp/tensor/_print.py b/dpnp/tensor/_print.py index 708efc8e2ad2..715a0969d2b5 100644 --- a/dpnp/tensor/_print.py +++ b/dpnp/tensor/_print.py @@ -95,17 +95,14 @@ def _options_dict( val = local[str_arg] if val is not None: if not isinstance(val, str): - raise TypeError( - "`{}` ".format(str_arg) + "must be of `string` type." - ) + raise TypeError(f"`{str_arg}` must be of `string` type.") options[str_arg] = val signs = ["-", "+", " "] if sign is not None: if sign not in signs: raise ValueError( - "`sign` must be one of" - + ", ".join("`{}`".format(s) for s in signs) + "`sign` must be one of" + ", ".join(f"`{s}`" for s in signs) ) options["sign"] = sign @@ -114,7 +111,7 @@ def _options_dict( if floatmode not in floatmodes: raise ValueError( "`floatmode` must be one of" - + ", ".join("`{}`".format(m) for m in floatmodes) + + ", ".join(f"`{m}`" for m in floatmodes) ) options["floatmode"] = floatmode diff --git a/dpnp/tensor/_slicing.pyx b/dpnp/tensor/_slicing.pyx index 10676e63397d..1d0f4afdd28c 100644 --- a/dpnp/tensor/_slicing.pyx +++ b/dpnp/tensor/_slicing.pyx @@ -208,8 +208,8 @@ def _basic_slice_meta(ind, shape : tuple, strides : tuple, offset : int): ) else: raise IndexError( - "Index {0} is out of range for axes 0 with " - "size {1}".format(ind, shape[0])) + f"Index {ind} is out of range for axes 0 with " + f"size {shape[0]}") elif isinstance(ind, (ndarray, usm_ndarray)): return (shape, strides, offset, (ind,), 0) elif isinstance(ind, tuple): @@ -273,8 +273,8 @@ def _basic_slice_meta(ind, shape : tuple, strides : tuple, offset : int): if axes_referenced > len(shape): raise IndexError( "too many indices for an array, array is " - "{0}-dimensional, but {1} were indexed".format( - len(shape), axes_referenced)) + f"{len(shape)}-dimensional, but " + f"{axes_referenced} were indexed") if ellipses_count: ellipses_count = len(shape) - axes_referenced new_shape_len = (newaxis_count + ellipses_count @@ -335,8 +335,8 @@ def _basic_slice_meta(ind, shape : tuple, strides : tuple, offset : int): 0 <= ind_i < shape[k] or -shape[k] <= ind_i < 0 ): raise IndexError( - "Index {0} is out of range for axes " - "{1} with size {2}".format(ind_i, k, shape[k]) + f"Index {ind_i} is out of range for axes " + f"{k} with size {shape[k]}" ) new_advanced_ind.append(ind_i) k_new = k + 1 @@ -359,8 +359,8 @@ def _basic_slice_meta(ind, shape : tuple, strides : tuple, offset : int): k = k_new else: raise IndexError( - "Index {0} is out of range for axes " - "{1} with size {2}".format(ind_i, k, shape[k]) + f"Index {ind_i} is out of range for axes " + f"{k} with size {shape[k]}" ) elif isinstance(ind_i, (ndarray, usm_ndarray)): if not array_streak: diff --git a/dpnp/tensor/libtensor/include/kernels/accumulators.hpp b/dpnp/tensor/libtensor/include/kernels/accumulators.hpp index 079fcf5e9c96..3c862c181818 100644 --- a/dpnp/tensor/libtensor/include/kernels/accumulators.hpp +++ b/dpnp/tensor/libtensor/include/kernels/accumulators.hpp @@ -74,9 +74,17 @@ struct NonZeroIndicator { static constexpr outputT out_one(1); static constexpr outputT out_zero(0); - static constexpr inputT val_zero(0); - return (val == val_zero) ? out_zero : out_one; + if constexpr (std::is_same_v) { + // NumPy treats any non-zero byte as True; read the raw byte + // rather than the bool value, see gh-2121 + const std::uint8_t u = sycl::bit_cast(val); + return (u == std::uint8_t{0}) ? out_zero : out_one; + } + else { + static constexpr inputT val_zero(0); + return (val == val_zero) ? out_zero : out_one; + } } }; @@ -85,7 +93,12 @@ struct NoOpTransformer { constexpr NoOpTransformer() {} - T operator()(const T &val) const { return val; } + // bool is normalized: a byte other than 0x00/0x01 would otherwise reach + // the scan operation and leak into the result, see gh-2121 + T operator()(const T &val) const + { + return dpnp::tensor::type_utils::normalize_bool(val); + } }; template @@ -920,10 +933,12 @@ sycl::event update_local_chunks(sycl::queue &exec_q, sycl::event dependent_event) { static constexpr NoOpIndexer out_indexer{}; - static constexpr NoOpIndexer iter_out_indexer{}; + // src rows stride by src_size; NoOp iter offset only ok for iter_nelems==1 + const Strided1DIndexer iter_out_indexer{/* size */ iter_nelems, + /* step */ src_size}; return final_update_local_chunks( + Strided1DIndexer, NoOpIndexer, ScanOpT>( exec_q, iter_nelems, src, src_size, local_scans, chunk_size, local_stride, iter_out_indexer, out_indexer, dependent_event); } @@ -1301,8 +1316,10 @@ struct Cumsum1DContigFactory { if constexpr (std::is_integral_v) { using cumsumT = std::int64_t; + // CastTransformer, not NoOpTransformer: an implicit bool + // conversion would read the raw byte, see gh-2121 fnT fn = - cumsum_val_contig_impl>; + cumsum_val_contig_impl>; return fn; } else { @@ -1419,8 +1436,10 @@ struct Cumsum1DStridedFactory { if constexpr (std::is_integral_v) { using cumsumT = std::int64_t; - fnT fn = - cumsum_val_strided_impl>; + // CastTransformer, not NoOpTransformer: an implicit bool + // conversion would read the raw byte, see gh-2121 + fnT fn = cumsum_val_strided_impl>; return fn; } else { diff --git a/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common.hpp b/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common.hpp index bb310272e7a6..39e472ddb727 100644 --- a/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common.hpp +++ b/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common.hpp @@ -46,6 +46,7 @@ #include "utils/offset_utils.hpp" #include "utils/sycl_alloc_utils.hpp" #include "utils/sycl_utils.hpp" +#include "utils/type_utils.hpp" #include "kernels/alignment.hpp" #include "kernels/dpnp_tensor_types.hpp" @@ -59,6 +60,7 @@ using dpnp::tensor::kernels::alignment_utils::required_alignment; using dpnp::tensor::sycl_utils::sub_group_load; using dpnp::tensor::sycl_utils::sub_group_store; +using dpnp::tensor::type_utils::normalize_bool; /*! @brief Functor for unary function evaluation on contiguous array */ template 1)) { + !std::is_same_v && (vec_sz > 1)) { auto sg = ndit.get_sub_group(); const std::uint16_t sgSize = sg.get_max_local_range()[0]; @@ -178,7 +182,7 @@ struct UnaryContigFunctor sub_group_load(sg, in_multi_ptr); #pragma unroll for (std::uint32_t k = 0; k < vec_sz; ++k) { - arg_vec[k] = op(arg_vec[k]); + arg_vec[k] = op(normalize_bool(arg_vec[k])); } sub_group_store(sg, arg_vec, out_multi_ptr); } @@ -186,7 +190,7 @@ struct UnaryContigFunctor else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < nelems_; k += sgSize) { - out[k] = op(in[k]); + out[k] = op(normalize_bool(in[k])); } } } @@ -216,7 +220,7 @@ struct UnaryContigFunctor sycl::vec res_vec; #pragma unroll for (std::uint8_t k = 0; k < vec_sz; ++k) { - res_vec[k] = op(arg_vec[k]); + res_vec[k] = op(normalize_bool(arg_vec[k])); } sub_group_store(sg, res_vec, out_multi_ptr); } @@ -224,7 +228,7 @@ struct UnaryContigFunctor else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < nelems_; k += sgSize) { - out[k] = op(in[k]); + out[k] = op(normalize_bool(in[k])); } } } @@ -238,7 +242,7 @@ struct UnaryContigFunctor (gid / sgSize) * (elems_per_sg - sgSize) + gid; const std::size_t end = std::min(nelems_, start + elems_per_sg); for (std::size_t offset = start; offset < end; offset += sgSize) { - out[offset] = op(in[offset]); + out[offset] = op(normalize_bool(in[offset])); } } } @@ -268,7 +272,7 @@ struct UnaryStridedFunctor UnaryOpT op{}; - res_[res_offset] = op(inp_[inp_offset]); + res_[res_offset] = op(normalize_bool(inp_[inp_offset])); } }; @@ -419,9 +423,13 @@ struct BinaryContigFunctor /* Each work-item processes vec_sz elements, contiguous in memory */ /* NOTE: work-group size must be divisible by sub-group size */ + // bool is excluded from the vector path: a byte other than 0x00/0x01 + // cannot be normalized element-wise there, see gh-2121 if constexpr (enable_sg_loadstore && BinaryOperatorT::supports_sg_loadstore::value && - BinaryOperatorT::supports_vec::value && (vec_sz > 1)) { + BinaryOperatorT::supports_vec::value && + !std::is_same_v && + !std::is_same_v && (vec_sz > 1)) { auto sg = ndit.get_sub_group(); std::uint16_t sgSize = sg.get_max_local_range()[0]; @@ -491,8 +499,8 @@ struct BinaryContigFunctor sycl::vec res_vec; #pragma unroll for (std::uint8_t vec_id = 0; vec_id < vec_sz; ++vec_id) { - res_vec[vec_id] = - op(arg1_vec[vec_id], arg2_vec[vec_id]); + res_vec[vec_id] = op(normalize_bool(arg1_vec[vec_id]), + normalize_bool(arg2_vec[vec_id])); } sub_group_store(sg, res_vec, out_multi_ptr); } @@ -500,7 +508,7 @@ struct BinaryContigFunctor else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < nelems_; k += sgSize) { - out[k] = op(in1[k], in2[k]); + out[k] = op(normalize_bool(in1[k]), normalize_bool(in2[k])); } } } @@ -514,7 +522,8 @@ struct BinaryContigFunctor (gid / sgSize) * (elems_per_sg - sgSize) + gid; const std::size_t end = std::min(nelems_, start + elems_per_sg); for (std::size_t offset = start; offset < end; offset += sgSize) { - out[offset] = op(in1[offset], in2[offset]); + out[offset] = op(normalize_bool(in1[offset]), + normalize_bool(in2[offset])); } } } @@ -553,7 +562,8 @@ struct BinaryStridedFunctor const auto &out_offset = three_offsets_.get_third_offset(); BinaryOperatorT op{}; - out[out_offset] = op(in1[inp1_offset], in2[inp2_offset]); + out[out_offset] = op(normalize_bool(in1[inp1_offset]), + normalize_bool(in2[inp2_offset])); } }; @@ -610,14 +620,15 @@ struct BinaryContigMatrixContigRowBroadcastingFunctor const argT1 mat_el = sub_group_load(sg, in1_multi_ptr); const argT2 vec_el = sub_group_load(sg, in2_multi_ptr); - resT res_el = op(mat_el, vec_el); + resT res_el = op(normalize_bool(mat_el), normalize_bool(vec_el)); sub_group_store(sg, res_el, out_multi_ptr); } else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < n_elems; k += sgSize) { - res[k] = op(mat[k], padded_vec[k % n1]); + res[k] = op(normalize_bool(mat[k]), + normalize_bool(padded_vec[k % n1])); } } } @@ -675,14 +686,15 @@ struct BinaryContigRowContigMatrixBroadcastingFunctor const argT2 mat_el = sub_group_load(sg, in2_multi_ptr); const argT1 vec_el = sub_group_load(sg, in1_multi_ptr); - resT res_el = op(vec_el, mat_el); + resT res_el = op(normalize_bool(vec_el), normalize_bool(mat_el)); sub_group_store(sg, res_el, out_multi_ptr); } else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < n_elems; k += sgSize) { - res[k] = op(padded_vec[k % n1], mat[k]); + res[k] = op(normalize_bool(padded_vec[k % n1]), + normalize_bool(mat[k])); } } } diff --git a/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp b/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp index 9384ec603754..fa5d3ce58d22 100644 --- a/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp +++ b/dpnp/tensor/libtensor/include/kernels/elementwise_functions/common_inplace.hpp @@ -44,6 +44,7 @@ #include "utils/offset_utils.hpp" #include "utils/sycl_alloc_utils.hpp" #include "utils/sycl_utils.hpp" +#include "utils/type_utils.hpp" #include "kernels/alignment.hpp" #include "kernels/dpnp_tensor_types.hpp" @@ -59,6 +60,22 @@ using dpnp::tensor::kernels::alignment_utils::required_alignment; using dpnp::tensor::sycl_utils::sub_group_load; using dpnp::tensor::sycl_utils::sub_group_store; +using dpnp::tensor::type_utils::normalize_bool; + +// a bool lhs is read as a value by the functors, so normalize it too and write +// the result back as a canonical byte, see gh-2121 +template +void apply_inplace(OpT &op, resT &res, const argT &rhs) +{ + if constexpr (std::is_same_v) { + resT tmp = normalize_bool(res); + op(tmp, normalize_bool(rhs)); + res = tmp; + } + else { + op(res, normalize_bool(rhs)); + } +} template 1)) { + !std::is_same_v && (vec_sz > 1)) { auto sg = ndit.get_sub_group(); std::uint16_t sgSize = sg.get_max_local_range()[0]; @@ -154,7 +173,7 @@ struct BinaryInplaceContigFunctor sub_group_load(sg, lhs_multi_ptr); #pragma unroll for (std::uint8_t vec_id = 0; vec_id < vec_sz; ++vec_id) { - op(res_vec[vec_id], arg_vec[vec_id]); + apply_inplace(op, res_vec[vec_id], arg_vec[vec_id]); } sub_group_store(sg, res_vec, lhs_multi_ptr); } @@ -162,7 +181,7 @@ struct BinaryInplaceContigFunctor else { const std::size_t lane_id = sg.get_local_id()[0]; for (std::size_t k = base + lane_id; k < nelems_; k += sgSize) { - op(lhs[k], rhs[k]); + apply_inplace(op, lhs[k], rhs[k]); } } } @@ -176,7 +195,7 @@ struct BinaryInplaceContigFunctor (gid / sgSize) * (elems_per_sg - sgSize) + gid; const std::size_t end = std::min(nelems_, start + elems_per_sg); for (std::size_t offset = start; offset < end; offset += sgSize) { - op(lhs[offset], rhs[offset]); + apply_inplace(op, lhs[offset], rhs[offset]); } } } @@ -210,7 +229,7 @@ struct BinaryInplaceStridedFunctor const auto &lhs_offset = two_offsets_.get_second_offset(); BinaryInplaceOperatorT op{}; - op(lhs[lhs_offset], rhs[inp_offset]); + apply_inplace(op, lhs[lhs_offset], rhs[inp_offset]); } }; @@ -257,14 +276,14 @@ struct BinaryInplaceRowMatrixBroadcastingFunctor const argT vec_el = sub_group_load(sg, in_multi_ptr); resT mat_el = sub_group_load(sg, out_multi_ptr); - op(mat_el, vec_el); + apply_inplace(op, mat_el, vec_el); sub_group_store(sg, mat_el, out_multi_ptr); } else { const std::size_t start = base + sg.get_local_id()[0]; for (std::size_t k = start; k < n_elems; k += sgSize) { - op(mat[k], padded_vec[k % n1]); + apply_inplace(op, mat[k], padded_vec[k % n1]); } } } diff --git a/dpnp/tensor/libtensor/include/kernels/reductions.hpp b/dpnp/tensor/libtensor/include/kernels/reductions.hpp index 42bea6f28126..9628e4028b26 100644 --- a/dpnp/tensor/libtensor/include/kernels/reductions.hpp +++ b/dpnp/tensor/libtensor/include/kernels/reductions.hpp @@ -56,6 +56,8 @@ namespace dpnp::tensor::kernels { +using dpnp::tensor::type_utils::normalize_bool; + using dpnp::tensor::ssize_t; namespace su_ns = dpnp::tensor::sycl_utils; @@ -1913,7 +1915,7 @@ struct SequentialSearchReduction const ssize_t inp_reduction_offset = inp_reduced_dims_indexer_(m); const ssize_t inp_offset = inp_iter_offset + inp_reduction_offset; - argT val = inp_[inp_offset]; + argT val = normalize_bool(inp_[inp_offset]); if (val == red_val) { idx_val = idx_reduction_op_(idx_val, static_cast(m)); } @@ -2058,7 +2060,7 @@ struct SearchReduction inp_reduced_dims_indexer_(arg_reduce_gid); auto inp_offset = inp_iter_offset + inp_reduction_offset; - argT val = inp_[inp_offset]; + argT val = normalize_bool(inp_[inp_offset]); if (val == local_red_val) { if constexpr (!First) { local_idx = @@ -2216,7 +2218,7 @@ struct CustomSearchReduction inp_reduced_dims_indexer_(arg_reduce_gid); auto inp_offset = inp_iter_offset + inp_reduction_offset; - argT val = inp_[inp_offset]; + argT val = normalize_bool(inp_[inp_offset]); if (val == local_red_val) { if constexpr (!First) { local_idx = diff --git a/dpnp/tensor/libtensor/include/kernels/sorting/isin.hpp b/dpnp/tensor/libtensor/include/kernels/sorting/isin.hpp index 2388f23fdf61..918eeb7e2f07 100644 --- a/dpnp/tensor/libtensor/include/kernels/sorting/isin.hpp +++ b/dpnp/tensor/libtensor/include/kernels/sorting/isin.hpp @@ -43,6 +43,7 @@ #include "kernels/sorting/search_sorted_detail.hpp" #include "utils/offset_utils.hpp" #include "utils/rich_comparisons.hpp" +#include "utils/type_utils.hpp" namespace dpnp::tensor::kernels { @@ -87,7 +88,10 @@ struct IsinFunctor static constexpr Compare comp{}; const std::size_t i = id[0]; - const T needle_v = needles_tp[needles_indexer(i)]; + // normalize: for bool a byte other than 0x00/0x01 would not compare + // equal to the normalized value in the hay array, see gh-2121 + const T needle_v = dpnp::tensor::type_utils::normalize_bool( + needles_tp[needles_indexer(i)]); // position of the needle_v in the hay array std::size_t pos{}; @@ -100,7 +104,10 @@ struct IsinFunctor // needle_v) is false, i.e. needle_v <= hay[pos] pos = search_sorted_detail::lower_bound_indexed_impl( hay_tp, zero, hay_nelems, needle_v, comp, hay_indexer); - bool out = (pos == hay_nelems ? false : hay_tp[pos] == needle_v); + bool out = + (pos == hay_nelems ? false + : dpnp::tensor::type_utils::normalize_bool( + hay_tp[pos]) == needle_v); out_tp[out_indexer(i)] = (invert) ? !out : out; } }; diff --git a/dpnp/tensor/libtensor/include/kernels/sorting/radix_sort.hpp b/dpnp/tensor/libtensor/include/kernels/sorting/radix_sort.hpp index 27eade9f3586..163f2ae64dcc 100644 --- a/dpnp/tensor/libtensor/include/kernels/sorting/radix_sort.hpp +++ b/dpnp/tensor/libtensor/include/kernels/sorting/radix_sort.hpp @@ -50,6 +50,7 @@ #include "kernels/dpnp_tensor_types.hpp" #include "kernels/sorting/sort_utils.hpp" #include "utils/sycl_alloc_utils.hpp" +#include "utils/type_utils.hpp" namespace dpnp::tensor::kernels { @@ -116,12 +117,15 @@ std::uint32_t ceil_log2(SizeT n) //---------------------------------------------------------- template -bool order_preserving_cast(bool val) +bool order_preserving_cast(const bool &val) { + // by reference: a bool copy lets the compiler assume a 0/1 byte, and the + // bucket index below reads only the low radix bits, see gh-2121 + const bool v = dpnp::tensor::type_utils::normalize_bool(val); if constexpr (is_ascending) - return val; + return v; else - return !val; + return !v; } template || std::is_same_v || std::is_same_v); +struct BoolLess +{ + bool operator()(const bool &v1, const bool &v2) const + { + using dpnp::tensor::type_utils::normalize_bool; + return !normalize_bool(v1) && normalize_bool(v2); + } +}; + +struct BoolGreater +{ + bool operator()(const bool &v1, const bool &v2) const + { + using dpnp::tensor::type_utils::normalize_bool; + return normalize_bool(v1) && !normalize_bool(v2); + } +}; + } // namespace detail template @@ -126,6 +146,12 @@ struct AscendingSorter std::less>; }; +template <> +struct AscendingSorter +{ + using type = detail::BoolLess; +}; + template struct AscendingSorter> { @@ -140,6 +166,12 @@ struct DescendingSorter std::greater>; }; +template <> +struct DescendingSorter +{ + using type = detail::BoolGreater; +}; + template struct DescendingSorter> { diff --git a/dpnp/tensor/libtensor/include/utils/type_utils.hpp b/dpnp/tensor/libtensor/include/utils/type_utils.hpp index b62bfb38a2e6..65d9e1a14b13 100644 --- a/dpnp/tensor/libtensor/include/utils/type_utils.hpp +++ b/dpnp/tensor/libtensor/include/utils/type_utils.hpp @@ -60,11 +60,28 @@ struct is_complex< template inline constexpr bool is_complex_v = is_complex::value; +// NumPy reads any non-zero bool byte as True; a non-canonical byte +// (not 0x00/0x01) used as a value would disagree, see gh-2121 +template +T normalize_bool(const T &v) +{ + if constexpr (std::is_same_v) { + // read the raw storage byte and test non-zero; using a + // non-canonical bool as a value is unreliable (UB) + const std::uint8_t u = sycl::bit_cast(v); + return u != std::uint8_t{0}; + } + else { + return v; + } +} + template dstTy convert_impl(const srcTy &v) { if constexpr (std::is_same_v) { - return v; + // bool needs normalizing even here, the byte may not be 0x00/0x01 + return normalize_bool(v); } else if constexpr (std::is_same_v) { if constexpr (is_complex_v) { diff --git a/dpnp/tests/tensor/test_tensor_accumulation.py b/dpnp/tests/tensor/test_tensor_accumulation.py index b7ea9147e100..d37cb706e3e5 100644 --- a/dpnp/tests/tensor/test_tensor_accumulation.py +++ b/dpnp/tests/tensor/test_tensor_accumulation.py @@ -126,6 +126,23 @@ def test_strided_cumsum_axis_sint(dt): assert dpt.all(res == dpt.expand_dims(expected, axis=1)) +@pytest.mark.parametrize("func", ["cumulative_sum", "cumulative_prod"]) +def test_batched_multilevel_scan(func): + # Regression test for gh-3063: multi-row scan over an axis needing >=3 + # levels (axis > chunk_size**2, chunk_size <= 2048) mis-strided rows. + get_queue_or_skip() + n0, n1 = 3, 5_000_000 + x = dpt.ones((n0, n1), dtype="i1") + + if func == "cumulative_prod": + x[:, 0] = 2 # leading 2 then ones -> cumprod is 2 everywhere + res = dpt.cumulative_prod(x, axis=1, dtype="i4") + assert dpt.all(res == 2) + else: + res = dpt.cumulative_sum(x, axis=1, dtype="i4") + assert dpt.all(res == dpt.arange(1, n1 + 1, dtype="i4")) + + def test_accumulate_scalar(): get_queue_or_skip() diff --git a/dpnp/tests/tensor/test_usm_ndarray_capi.py b/dpnp/tests/tensor/test_usm_ndarray_capi.py index d0e44f6712da..479c288698a5 100644 --- a/dpnp/tests/tensor/test_usm_ndarray_capi.py +++ b/dpnp/tests/tensor/test_usm_ndarray_capi.py @@ -51,7 +51,7 @@ def _pyx_capi_fnptr_to_callable( cap = mod.__pyx_capi__.get(pyx_capi_name, None) if cap is None: raise ValueError( - "__pyx_capi__ does not export {} capsule".format(pyx_capi_name) + f"__pyx_capi__ does not export {pyx_capi_name} capsule" ) # construct Python callable to invoke these functions cap_ptr_fn = ctypes.pythonapi.PyCapsule_GetPointer @@ -533,7 +533,7 @@ def _pyx_capi_int(X, pyx_capi_name, caps_name=b"int", val_restype=ctypes.c_int): cap = mod.__pyx_capi__.get(pyx_capi_name, None) if cap is None: raise ValueError( - "__pyx_capi__ does not export {} capsule".format(pyx_capi_name) + f"__pyx_capi__ does not export {pyx_capi_name} capsule" ) # construct Python callable to invoke these functions cap_ptr_fn = ctypes.pythonapi.PyCapsule_GetPointer diff --git a/dpnp/tests/tensor/test_usm_ndarray_ctor.py b/dpnp/tests/tensor/test_usm_ndarray_ctor.py index 95dd6469d1c5..4e2f46d64935 100644 --- a/dpnp/tests/tensor/test_usm_ndarray_ctor.py +++ b/dpnp/tests/tensor/test_usm_ndarray_ctor.py @@ -416,7 +416,7 @@ def test_slice_constructor_1d(): ]: assert np.array_equal( dpt.asnumpy(Xusm[ind]), Xh[ind] - ), "Failed for {}".format(ind) + ), f"Failed for {ind}" def test_slice_constructor_3d(): @@ -438,7 +438,7 @@ def test_slice_constructor_3d(): ]: assert np.array_equal( dpt.to_numpy(Xusm[ind]), Xh[ind] - ), "Failed for {}".format(ind) + ), f"Failed for {ind}" @pytest.mark.parametrize("usm_type", ["device", "shared", "host"]) @@ -451,7 +451,7 @@ def test_slice_suai(usm_type): for ind in [slice(2, 3, None), slice(5, 7, None), slice(3, 9, None)]: assert np.array_equal( dpm.as_usm_memory(Xusm[ind]).copy_to_host(), Xh[ind] - ), "Failed for {}".format(ind) + ), f"Failed for {ind}" def test_slicing_basic(): diff --git a/dpnp/tests/tensor/test_usm_ndarray_manipulation.py b/dpnp/tests/tensor/test_usm_ndarray_manipulation.py index bf0bd3226309..cd5962615fb9 100644 --- a/dpnp/tests/tensor/test_usm_ndarray_manipulation.py +++ b/dpnp/tests/tensor/test_usm_ndarray_manipulation.py @@ -34,6 +34,7 @@ from numpy.testing import assert_, assert_array_equal, assert_raises_regex import dpnp.tensor as dpt +import dpnp.tensor._tensor_impl as ti from dpnp.tensor._numpy_helper import AxisError from .helper import get_queue_or_skip @@ -1412,6 +1413,24 @@ def test_repeat_strided_repeats(): assert dpt.all(res == x) +def test_repeat_nonstandard_bool_bytes(): + # NumPy treats any non-zero byte of a bool as True, see gh-2121 + get_queue_or_skip() + + raw = dpt.asarray([0, 1, 2, 255, 0, 1], dtype="u1") + reps = dpt.usm_ndarray(raw.shape, dtype="?", buffer=raw.usm_data) + x = dpt.arange(reps.size, dtype="i4") + + res = dpt.repeat(x, reps) + assert_array_equal(dpt.asnumpy(res), np.array([1, 2, 3, 5], dtype="i4")) + + # `repeat` casts a non-int64 `reps` first, so drive the scan directly too + cumsum = dpt.empty(reps.size, dtype="i8") + total = ti._cumsum_1d(reps, cumsum, sycl_queue=reps.sycl_queue) + assert total == 4 + assert_array_equal(dpt.asnumpy(cumsum), np.array([0, 1, 2, 3, 3, 4])) + + def test_repeat_size1_repeats(): get_queue_or_skip() diff --git a/dpnp/tests/test_indexing.py b/dpnp/tests/test_indexing.py index 84bf62d03562..378ca2fa10e6 100644 --- a/dpnp/tests/test_indexing.py +++ b/dpnp/tests/test_indexing.py @@ -277,6 +277,29 @@ def test_place_insert_from_empty_vals(self, xp): def test_place_wrong_array_type(self, xp): assert_raises(TypeError, xp.place, [1, 2, 3], [True, False], [0, 1]) + # NumPy treats any non-zero byte of a bool as True, see gh-2121 + def test_extract_nonstandard_bool_bytes(self): + raw = numpy.array([0, 1, 2, 255, 0, 1], dtype=numpy.uint8) + a = numpy.arange(raw.size, dtype=numpy.int32) + mask = raw.view(numpy.bool_) + ia = dpnp.asarray(a) + imask = dpnp.asarray(raw).view(dpnp.bool) + + result = dpnp.extract(imask, ia) + expected = numpy.extract(mask, a) + assert_array_equal(result, expected) + + def test_place_nonstandard_bool_bytes(self): + raw = numpy.array([0, 1, 2, 255, 0, 1], dtype=numpy.uint8) + a = numpy.arange(raw.size, dtype=numpy.int32) + mask = raw.view(numpy.bool_) + ia = dpnp.asarray(a) + imask = dpnp.asarray(raw).view(dpnp.bool) + + dpnp.place(ia, imask, [-1]) + numpy.place(a, mask, [-1]) + assert_array_equal(ia, a) + @pytest.mark.parametrize("dt", get_all_dtypes(no_none=True)) def test_both(self, dt): a = numpy.random.rand(10).astype(dt) @@ -616,6 +639,42 @@ def test_array_method(self, dtype): ia = dpnp.array(a) assert_array_equal(a.nonzero(), ia.nonzero()) + # NumPy treats any non-zero byte of a bool as True, see gh-2121 + @pytest.mark.parametrize( + "bytes_val", + [ + [0, 1, 2, 255, 0, 1], + [2] * 8, + [255], + [0] * 8, + [0, 128] * 64, + list(range(256)), + ], + ids=["mixed", "all_twos", "single_255", "all_zeros", "long", "range"], + ) + def test_nonstandard_bool_bytes(self, bytes_val): + a = numpy.array(bytes_val, dtype=numpy.uint8).view(numpy.bool_) + ia = dpnp.asarray(numpy.array(bytes_val, dtype=numpy.uint8)).view( + dpnp.bool + ) + + assert_array_equal(numpy.nonzero(a), dpnp.nonzero(ia)) + assert_array_equal(numpy.where(a), dpnp.where(ia)) + + def test_nonstandard_bool_bytes_strided(self): + raw = numpy.arange(24, dtype=numpy.uint8) + a = raw.view(numpy.bool_)[::3] + ia = dpnp.asarray(raw).view(dpnp.bool)[::3] + + assert_array_equal(numpy.nonzero(a), dpnp.nonzero(ia)) + + def test_nonstandard_bool_bytes_2d(self): + raw = numpy.array([[0, 1, 2], [255, 0, 7]], dtype=numpy.uint8) + a = raw.view(numpy.bool_) + ia = dpnp.asarray(raw).view(dpnp.bool) + + assert_array_equal(numpy.nonzero(a), dpnp.nonzero(ia)) + class TestPut: @pytest.mark.parametrize("a_dt", get_all_dtypes(no_none=True)) diff --git a/dpnp/tests/test_linalg.py b/dpnp/tests/test_linalg.py index 9cd9a1b9b8fe..fe00f4211421 100644 --- a/dpnp/tests/test_linalg.py +++ b/dpnp/tests/test_linalg.py @@ -1690,6 +1690,163 @@ def test_path(self): assert expected[0] == result[0] assert expected[1] == result[1] + @pytest.mark.parametrize( + "subscripts, shape1, shape2", + [ + ("lkz,lxpq->kxpqz", (3, 2, 2), (3, 1, 6, 6)), + ("lkz,lxpq->kxpqz", (4, 3, 2), (4, 2, 5, 5)), + ("ij,jk->ik", (4, 5), (5, 6)), + ("ijk,ikl->ijl", (2, 3, 4), (2, 4, 5)), + ("lk,lpq->kpq", (3, 2), (3, 6, 6)), + ], + ) + def test_contraction_order_k(self, subscripts, shape1, shape2): + # for order="K" (the default), a contraction is materialized into a + # newly allocated array, so the result is c-contiguous when the + # operands are, matching NumPy + a = generate_random_numpy_array(shape1) + b = generate_random_numpy_array(shape2) + ia, ib = dpnp.array(a), dpnp.array(b) + + result = dpnp.einsum(subscripts, ia, ib) + expected = numpy.einsum(subscripts, a, b) + assert result.flags.c_contiguous == expected.flags.c_contiguous + assert result.flags.c_contiguous + assert_dtype_allclose(result, expected) + + @pytest.mark.parametrize("optimize", [False, True, "greedy", "optimal"]) + @pytest.mark.parametrize("order", ["C", "F", "A", "K", None]) + @pytest.mark.parametrize("order1", ["C", "F"]) + @pytest.mark.parametrize("order2", ["C", "F"]) + def test_contraction_order(self, optimize, order, order1, order2): + if order is None and optimize is not False: + pytest.skip("numpy raises AttributeError for order=None here") + a = generate_random_numpy_array((4, 5), order=order1) + b = generate_random_numpy_array((5, 6), order=order2) + ia, ib = dpnp.array(a), dpnp.array(b) + + result = dpnp.einsum( + "ij,jk->ik", ia, ib, order=order, optimize=optimize + ) + expected = numpy.einsum( + "ij,jk->ik", a, b, order=order, optimize=optimize + ) + assert result.flags.c_contiguous == expected.flags.c_contiguous + assert result.flags.f_contiguous == expected.flags.f_contiguous + assert_dtype_allclose(result, expected) + + @pytest.mark.parametrize("optimize", [True, "greedy", "optimal"]) + @pytest.mark.parametrize( + "subscripts, shapes", + [ + ("...ft,mf->...mt", [(2, 3, 5), (4, 3)]), + ("lk,lpq->kpq", [(3, 4), (3, 5, 6)]), + ("ijk,jkl->il", [(2, 3, 4), (3, 4, 5)]), + ("ij,jk,kl->il", [(4, 5), (5, 6), (6, 7)]), + ], + ) + @pytest.mark.parametrize("order1", ["C", "F"]) + @pytest.mark.parametrize("order2", ["C", "F"]) + def test_contraction_order_k_optimize( + self, optimize, subscripts, shapes, order1, order2 + ): + # NumPy only copies the result into a c-contiguous array on its + # unoptimized path; an optimized one is matmul-based, as dpnp always + # is, and keeps the permuted layout that the contraction produces + orders = [order1, order2] + ["C"] * (len(shapes) - 2) + arrays = [ + generate_random_numpy_array(shape, order=o) + for shape, o in zip(shapes, orders) + ] + iarrays = [dpnp.array(a, order=o) for a, o in zip(arrays, orders)] + + result = dpnp.einsum(subscripts, *iarrays, optimize=optimize) + expected = numpy.einsum(subscripts, *arrays, optimize=optimize) + assert result.flags.c_contiguous == expected.flags.c_contiguous + assert result.flags.f_contiguous == expected.flags.f_contiguous + assert_dtype_allclose(result, expected) + + def test_contraction_order_a_trivial(self): + # an operand that is both c- and f-contiguous (here 1-D) is + # f_contiguous, so order="A" resolves to "F" as it does in NumPy + a = generate_random_numpy_array(4) + b = generate_random_numpy_array((4, 5, 6), order="F") + ia, ib = dpnp.array(a), dpnp.array(b, order="F") + + result = dpnp.einsum("i,ijk->jk", ia, ib, order="A") + expected = numpy.einsum("i,ijk->jk", a, b, order="A") + assert result.flags.c_contiguous == expected.flags.c_contiguous + assert result.flags.f_contiguous == expected.flags.f_contiguous + assert_dtype_allclose(result, expected) + + @pytest.mark.parametrize("order", ["C", "F", "A", "K", None]) + def test_empty_operand_order(self, order): + # a contraction over a size-0 dimension is all zeros, and `order` is + # honored for it as it is for a non-empty one + a = numpy.ones((2, 0)) + b = numpy.ones((0, 4)) + ia, ib = dpnp.array(a), dpnp.array(b) + + result = dpnp.einsum("ij,jk->ik", ia, ib, order=order) + expected = numpy.einsum("ij,jk->ik", a, b, order=order) + assert result.flags.c_contiguous == expected.flags.c_contiguous + assert result.flags.f_contiguous == expected.flags.f_contiguous + assert_dtype_allclose(result, expected) + + def test_empty_operand_out(self): + # `out` is filled with zeros and returned for a size-0 contraction + a = numpy.ones((2, 0)) + b = numpy.ones((0, 4)) + ia, ib = dpnp.array(a), dpnp.array(b) + iout = dpnp.full((2, 4), 9.0) + out = numpy.full((2, 4), 9.0) + + result = dpnp.einsum("ij,jk->ik", ia, ib, out=iout) + expected = numpy.einsum("ij,jk->ik", a, b, out=out) + assert result is iout + assert_dtype_allclose(result, expected) + + @pytest.mark.parametrize("subscripts", ["ij->ji", "ij->ij", "ii->i"]) + @pytest.mark.parametrize("order", ["C", "F", "A", "K", None]) + def test_unary_view_order(self, subscripts, order): + # a single-operand einsum with no summed index returns a view of the + # operand for every value of `order`, as it does in NumPy + # the dtype is pinned because the strides below scale with itemsize + a = generate_random_numpy_array((4, 4), dtype=dpnp.default_float_type()) + ia = dpnp.array(a) + + result = dpnp.einsum(subscripts, ia, order=order) + expected = numpy.einsum(subscripts, a, order=order) + assert result.get_array()._pointer == ia.get_array()._pointer + assert result.strides == expected.strides + assert_dtype_allclose(result, expected) + + @pytest.mark.parametrize("subscripts", ["ij->ji", "ii->i"]) + @pytest.mark.parametrize("order", ["C", "F", "A", "K"]) + def test_unary_view_is_writeable(self, subscripts, order): + # the view returned for a unary einsum without summation is writeable, + # so an assignment through it is visible in the operand + a = generate_random_numpy_array((4, 4)) + ia = dpnp.array(a) + + result = dpnp.einsum(subscripts, ia, order=order) + result[...] = 0 + expected = numpy.einsum(subscripts, a, order=order) + expected[...] = 0 + assert_dtype_allclose(ia, a) + + @pytest.mark.parametrize("order", ["W", "w", "", "CF"]) + def test_order_error(self, order): + a = dpnp.ones((3, 3)) + # a unary einsum without summation returns a view without going + # through dpnp.asarray, so `order` is validated up front + assert_raises(ValueError, dpnp.einsum, "ii->i", a, order=order) + assert_raises(ValueError, dpnp.einsum, "ij,jk->ik", a, a, order=order) + + def test_order_type_error(self): + a = dpnp.ones((3, 3)) + assert_raises(TypeError, dpnp.einsum, "ii->i", a, order=1) + class TestInv: @pytest.mark.parametrize( diff --git a/dpnp/tests/test_logic.py b/dpnp/tests/test_logic.py index a7696fe9852d..040bf0b00df7 100644 --- a/dpnp/tests/test_logic.py +++ b/dpnp/tests/test_logic.py @@ -1,3 +1,5 @@ +import operator + import dpctl import numpy import pytest @@ -900,3 +902,158 @@ def test_isin_errors(self): with pytest.raises(ExecutionPlacementError): dpnp.isin(a, b) + + +# NumPy treats any non-zero byte of a bool as True, see gh-2121 +class TestNonstandardBoolBytes: + def _views(self, values): + raw = numpy.array(values, dtype=numpy.uint8) + return raw.view(numpy.bool_), dpnp.asarray(raw).view(dpnp.bool) + + @pytest.mark.parametrize( + "op", + [ + "equal", + "not_equal", + "less", + "less_equal", + "greater", + "greater_equal", + "logical_and", + "logical_or", + "logical_xor", + "bitwise_and", + "bitwise_or", + "bitwise_xor", + "maximum", + "minimum", + ], + ) + def test_binary(self, op): + a, ia = self._views([0, 1, 2, 255, 3, 0]) + b, ib = self._views([1, 1, 0, 2, 0, 7]) + + result = getattr(dpnp, op)(ia, ib) + expected = getattr(numpy, op)(a, b) + assert_array_equal(result, expected) + + @pytest.mark.parametrize("op", ["iand", "ior", "ixor", "iadd", "imul"]) + @pytest.mark.parametrize("strided", [False, True]) + def test_inplace(self, op, strided): + # the destination is read as a value too, so its byte must normalize + a, ia = self._views([0, 1, 2, 255, 3, 0, 128, 127]) + b, ib = self._views([1, 1, 0, 2, 0, 7, 255, 1]) + if strided: + a, ia, b, ib = a[::2], ia[::2], b[::2], ib[::2] + + getattr(operator, op)(ia, ib) + getattr(operator, op)(a, b) + assert_array_equal(ia, a) + + raw = dpnp.asnumpy(ia).view(numpy.uint8) + assert numpy.all((raw == 0) | (raw == 1)) + + @pytest.mark.parametrize("op", ["logical_not", "bitwise_invert"]) + def test_unary(self, op): + a, ia = self._views([0, 1, 2, 255, 3, 0]) + + result = getattr(dpnp, op)(ia) + expected = getattr(numpy, op)(a) + assert_array_equal(result, expected) + + @pytest.mark.parametrize( + "op", ["sum", "prod", "max", "min", "all", "any", "argmax", "argmin"] + ) + def test_reduction(self, op): + a, ia = self._views([0, 1, 2, 255, 3, 0]) + + result = getattr(dpnp, op)(ia) + expected = getattr(numpy, op)(a) + assert_array_equal(result, expected) + + @pytest.mark.parametrize("op", ["cumulative_sum", "cumulative_prod"]) + @pytest.mark.parametrize("include_initial", [False, True]) + def test_cumulative_bool_out(self, op, include_initial): + # a bool accumulator must scan normalized values, not the raw bytes + a, ia = self._views([0, 1, 2, 255, 3, 0, 128, 127]) + + result = getattr(dpnp, op)( + ia, dtype=dpnp.bool, include_initial=include_initial + ) + expected = getattr(numpy, op)( + a, dtype=numpy.bool_, include_initial=include_initial + ) + assert_array_equal(result, expected) + + raw = dpnp.asnumpy(result).view(numpy.uint8) + assert numpy.all((raw == 0) | (raw == 1)) + + def test_result_has_canonical_bytes(self): + # a computed bool result must only ever hold 0x00 or 0x01; `sort` and + # `unique` are excluded, they select elements and so carry the bytes + _, ia = self._views([0, 1, 2, 255, 3, 0]) + _, ib = self._views([1, 1, 0, 2, 0, 7]) + + for res in ( + dpnp.equal(ia, ib), + dpnp.not_equal(ia, ib), + dpnp.logical_not(ia), + dpnp.logical_and(ia, ib), + dpnp.logical_or(ia, ib), + dpnp.logical_xor(ia, ib), + dpnp.bitwise_and(ia, ib), + dpnp.bitwise_or(ia, ib), + dpnp.bitwise_xor(ia, ib), + dpnp.maximum(ia, ib), + dpnp.isin(ia, ib), + dpnp.max(ia), + dpnp.min(ia), + dpnp.all(ia), + dpnp.any(ia), + ): + raw = dpnp.asnumpy(res).view(numpy.uint8) + assert numpy.all((raw == 0) | (raw == 1)) + + @pytest.mark.parametrize( + "values", + [[0, 1, 2, 255, 3, 0], [2] * 6, [255] * 6, [0, 128, 0, 127, 254, 1]], + ) + @pytest.mark.parametrize("kind", [None, "mergesort", "radixsort"]) + def test_sort_kinds_logically_ordered(self, values, kind): + # dpnp normalizes, so False elements sort before True ones; NumPy + # leaves the raw bytes in place, so only the logical order matches + a, ia = self._views(values) + kwargs = {} if kind is None else {"kind": kind} + + result = dpnp.asnumpy(dpnp.sort(ia, **kwargs)).view(numpy.uint8) != 0 + expected = numpy.sort(a).view(numpy.uint8) != 0 + assert_array_equal(result, expected) + + def test_unique(self): + a, ia = self._views([0, 1, 2, 255, 3, 0]) + + assert_array_equal(dpnp.unique(ia), numpy.unique(a)) + + def test_isin(self): + a, ia = self._views([0, 1, 2, 255, 3, 0]) + b, ib = self._views([1, 1, 0, 2, 0, 7]) + + assert_array_equal(dpnp.isin(ia, ib), numpy.isin(a, b)) + + def test_searchsorted(self): + # dpnp normalizes, so the needles compare as True; NumPy orders bool + # by raw byte, so only the normalized comparison matches + a, ia = self._views([0, 0, 1, 1]) + v, iv = self._views([2, 255, 0]) + + expected = numpy.searchsorted( + a.view(numpy.uint8) != 0, v.view(numpy.uint8) != 0 + ) + assert_array_equal(dpnp.searchsorted(ia, iv), expected) + + def test_argsort_is_logically_ordered(self): + raw = numpy.array([0, 1, 2, 255, 3, 0], dtype=numpy.uint8) + ia = dpnp.asarray(raw).view(dpnp.bool) + + gathered = raw[dpnp.asnumpy(dpnp.argsort(ia))] != 0 + assert_array_equal(gathered, numpy.sort(gathered))