From 2000d877d71e916ec0e0f6da835b287a43efe629 Mon Sep 17 00:00:00 2001 From: Anton <100830759+antonwolfy@users.noreply.github.com> Date: Mon, 7 Sep 2026 13:52:38 +0200 Subject: [PATCH 1/9] Remove obsolete array API skip for test_dunder_dlpack (#3061) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The array API conformance job skipped `test_dlpack.py::test_dunder_dlpack` as a workaround for a bug in the test itself (tracked in [array-api-tests#457](https://github.com/data-apis/array-api-tests/pull/457)): the test pinned the requested `dl_device` to `kDLCPU` while drawing `copy` from `{True, False, None}`, so on a SYCL device (reported as `kDLOneAPI`) the `copy=False` case forced a cross-device transfer and the spec-mandated `BufferError` was counted as a failure. The upstream test has since been fixed to tolerate `BufferError` only when `copy is False` and the requested device differs from the array's own `__dlpack_device__()` — behavior dpnp already implements correctly by delegating to dpctl. The conformance job checks out the array-api-tests default branch without pinning a ref, so the fix is already picked up on new runs and the skip is now dead weight. This change removes only that entry. The remaining `tanh` special-case skip is unrelated (an open array-api spec issue) and is left in place. --- .github/workflows/array-api-skips.txt | 3 --- 1 file changed, 3 deletions(-) diff --git a/.github/workflows/array-api-skips.txt b/.github/workflows/array-api-skips.txt index 2c94f9180fd..fb81a683e6d 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] From cbd45dff229ccebcea1e54e01827b5bb65c3efd4 Mon Sep 17 00:00:00 2001 From: Anton <100830759+antonwolfy@users.noreply.github.com> Date: Tue, 8 Sep 2026 19:31:01 +0200 Subject: [PATCH 2/9] Reuse reduction/dot buffer as sqrt output in linalg.norm (#3062) In the `dpnp.linalg.norm` implementation, the 2-norm and Frobenius-norm branches computed `sqrt(sum(...))` (or `sqrt(dot(...))` on the `axis=None` fast path) allocating a fresh array for the `sqrt` output on top of the array already produced by the reduction. These branches now capture the reduction (or dot) result and pass it back as `out=` to `dpnp.sqrt`, so the intermediate buffer is reused as the output instead of allocating a new one. The reused buffer is a private, unaliased array in every case, and its dtype matches the `sqrt` output, so the in-place write is safe. On the `axis=None` fast path the reused buffer is a scalar, so the change there is for consistency. This is an allocation saving only; it does not change results or reduce the number of kernel launches. --- CHANGELOG.md | 1 + dpnp/linalg/dpnp_utils_linalg.py | 9 +++++---- 2 files changed, 6 insertions(+), 4 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 09961c7ccd3..24064fcc16a 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -47,6 +47,7 @@ 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) ### Deprecated diff --git a/dpnp/linalg/dpnp_utils_linalg.py b/dpnp/linalg/dpnp_utils_linalg.py index 527235496b9..ac1c87a61b5 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 From d0548380bbc6deaec2dd7c3b25cd8215235e8e8e Mon Sep 17 00:00:00 2001 From: Anton <100830759+antonwolfy@users.noreply.github.com> Date: Wed, 9 Sep 2026 13:32:21 +0200 Subject: [PATCH 3/9] Fix row corruption in batched multi-level cumulative scan (#3063) ## Summary `dpnp.cumsum`, `dpnp.cumprod`, and their `nan`/`cumulative_*` variants (including `dpnp.tensor.cumulative_sum`/`cumulative_prod`) could silently return incorrect results when accumulating along an axis of a multi-row array, with no error raised. The corruption is triggered when the accumulation axis is long enough to require three or more reduction levels in the internal block-scan (axis length greater than `chunk_size ** 2`, where `chunk_size = wg_size * n_wi`, i.e. roughly a million elements on GPU and about four million on CPU) and the array has more than one batch row. The elementwise total is conserved but redistributed across rows. ## Root cause In the batched scan driver `inclusive_scan_iter` (`accumulators.hpp`), the intermediate block-scan fixup `update_local_chunks` passed a `NoOpIndexer` for the per-row (iter) offset into the `src` buffer. The intermediate block-scan buffers, however, are laid out strided per row with a stride of `src_size`. The `local_scans` read already used the correct `iter_gid * local_stride` offset, but the `src` write offset was unstrided, so rows overwrote each other. This cancels out only when `iter_nelems == 1` (a single row) or when fewer than three levels are needed, which is why the bug went unnoticed. ## Fix Give `update_local_chunks` a `Strided1DIndexer{iter_nelems, src_size}` for the row offset instead of the `NoOpIndexer` (the within-row `out_indexer` stays `NoOp`, since the level buffers are contiguous within a row). --- CHANGELOG.md | 1 + .../libtensor/include/kernels/accumulators.hpp | 6 ++++-- dpnp/tests/tensor/test_tensor_accumulation.py | 17 +++++++++++++++++ 3 files changed, 22 insertions(+), 2 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 24064fcc16a..1ddd33b3008 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -96,6 +96,7 @@ 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) ### Security diff --git a/dpnp/tensor/libtensor/include/kernels/accumulators.hpp b/dpnp/tensor/libtensor/include/kernels/accumulators.hpp index 079fcf5e9c9..fe3cec59a17 100644 --- a/dpnp/tensor/libtensor/include/kernels/accumulators.hpp +++ b/dpnp/tensor/libtensor/include/kernels/accumulators.hpp @@ -920,10 +920,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); } diff --git a/dpnp/tests/tensor/test_tensor_accumulation.py b/dpnp/tests/tensor/test_tensor_accumulation.py index b7ea9147e10..d37cb706e3e 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() From af5995c609881b36ba1f51a67576990448af20c2 Mon Sep 17 00:00:00 2001 From: Anton <100830759+antonwolfy@users.noreply.github.com> Date: Wed, 9 Sep 2026 15:49:55 +0200 Subject: [PATCH 4/9] Fix documentation rendering for dpnp.ndarray operator methods and dpnp.histogram2d (#3064) Fixes several documentation rendering issues. - `dpnp.ndarray` operator methods documented their behavior with a `:math:` role wrapping the expression in `\text{...}`. In LaTeX/MathJax math mode `&`, `^`, and `%` are special, so `__and__`, `__xor__`, `__mod__` and their variants rendered as error strings or truncated text. These are plain Python operator snippets, so they now render as inline code literals; all operator dunder docstrings were converted for consistency. - `dpnp.histogram2d` and `dpnp.left_shift` were missing the blank line before their `Returns` section, so numpydoc did not parse it as a section header. --- dpnp/dpnp_array.py | 120 +++++++++++++++++----------------- dpnp/dpnp_iface_bitwise.py | 1 + dpnp/dpnp_iface_histograms.py | 1 + 3 files changed, 62 insertions(+), 60 deletions(-) diff --git a/dpnp/dpnp_array.py b/dpnp/dpnp_array.py index 86055a4828f..b225fb2c732 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_iface_bitwise.py b/dpnp/dpnp_iface_bitwise.py index 604fd365ee1..350cc740782 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 8cd8bd3c08e..53a4239e7a1 100644 --- a/dpnp/dpnp_iface_histograms.py +++ b/dpnp/dpnp_iface_histograms.py @@ -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) From 039b7f14c865792b55ca91842165dbe2264ac163 Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty <59661409+abagusetty@users.noreply.github.com> Date: Wed, 9 Sep 2026 13:58:06 -0500 Subject: [PATCH 5/9] Fix `dpnp.einsum` memory-layout (#3058) One more minor layout issue detected during an app-testing in comparison to numpy behavior: Fixes: #3056 --- CHANGELOG.md | 1 + dpnp/dpnp_utils/dpnp_utils_einsum.py | 31 +++++- dpnp/tests/test_linalg.py | 157 +++++++++++++++++++++++++++ 3 files changed, 185 insertions(+), 4 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 1ddd33b3008..d55b09e4f92 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -97,6 +97,7 @@ This release is compatible with NumPy 2.5. * 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) ### Security diff --git a/dpnp/dpnp_utils/dpnp_utils_einsum.py b/dpnp/dpnp_utils/dpnp_utils_einsum.py index 842596bdc32..efa669126a1 100644 --- a/dpnp/dpnp_utils/dpnp_utils_einsum.py +++ b/dpnp/dpnp_utils/dpnp_utils_einsum.py @@ -1039,8 +1039,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 +1124,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. @@ -1226,6 +1243,12 @@ 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: + 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" + 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/tests/test_linalg.py b/dpnp/tests/test_linalg.py index 9cd9a1b9b8f..fe00f421142 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( From f418987ed569cb21cddf1a5e75371e47323dccfa Mon Sep 17 00:00:00 2001 From: Abhishek Bagusetty <59661409+abagusetty@users.noreply.github.com> Date: Thu, 10 Sep 2026 05:49:52 -0500 Subject: [PATCH 6/9] Fix nonzero bool nonstandard bytes (#3055) Boolean arrays may contain non-zero bytes like `0x02` or `0xFF` that NumPy treats as True but dpnp does not. Fixes #3054 --- CHANGELOG.md | 1 + .../include/kernels/accumulators.hpp | 29 +++- .../kernels/elementwise_functions/common.hpp | 46 +++-- .../elementwise_functions/common_inplace.hpp | 33 +++- .../libtensor/include/kernels/reductions.hpp | 8 +- .../include/kernels/sorting/isin.hpp | 11 +- .../include/kernels/sorting/radix_sort.hpp | 18 +- .../include/utils/rich_comparisons.hpp | 32 ++++ .../libtensor/include/utils/type_utils.hpp | 19 ++- .../tensor/test_usm_ndarray_manipulation.py | 19 +++ dpnp/tests/test_indexing.py | 59 +++++++ dpnp/tests/test_logic.py | 157 ++++++++++++++++++ 12 files changed, 392 insertions(+), 40 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index d55b09e4f92..0f47c903d54 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -98,6 +98,7 @@ This release is compatible with NumPy 2.5. * 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/dpnp/tensor/libtensor/include/kernels/accumulators.hpp b/dpnp/tensor/libtensor/include/kernels/accumulators.hpp index fe3cec59a17..3c862c18181 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 @@ -1303,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 { @@ -1421,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 bb310272e7a..39e472ddb72 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 9384ec60375..fa5d3ce58d2 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 42bea6f2812..9628e4028b2 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 2388f23fdf6..918eeb7e2f0 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 27eade9f358..163f2ae64dc 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 b62bfb38a2e..65d9e1a14b1 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_usm_ndarray_manipulation.py b/dpnp/tests/tensor/test_usm_ndarray_manipulation.py index bf0bd322630..cd5962615fb 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 84bf62d0356..378ca2fa10e 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_logic.py b/dpnp/tests/test_logic.py index a7696fe9852..040bf0b00df 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)) From 1d013525cd41218a6f71e7b921aafb61f12a670b Mon Sep 17 00:00:00 2001 From: Anton <100830759+antonwolfy@users.noreply.github.com> Date: Thu, 10 Sep 2026 14:43:18 +0200 Subject: [PATCH 7/9] Fix malformed :obj: cross-reference roles in docstrings (#3065) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two `See Also` entries used `obj:` instead of `:obj:` (missing the leading colon), so Sphinx did not recognize the cross-reference role and rendered it as literal `obj:` text in the API docs: - `dpnp.argsort` — the `dpnp.take_along_axis` entry rendered as "…from obj: dpnp.argsort…" - `dpnp.histogram` — the `dpnp.histogram_bin_edges` entry rendered as "…by the obj:dpnp.histogram function" Adding the missing colon makes both render as proper cross-reference links. This is a docstring-only change with no functional impact. --- dpnp/dpnp_iface_histograms.py | 2 +- dpnp/dpnp_iface_sorting.py | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/dpnp/dpnp_iface_histograms.py b/dpnp/dpnp_iface_histograms.py index 53a4239e7a1..1ac51d18f4c 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 -------- diff --git a/dpnp/dpnp_iface_sorting.py b/dpnp/dpnp_iface_sorting.py index 8f6f3e80f0d..5feffffbc43 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 From 53ac592c4cc6d3271394771c7d71e1b4ef81ae2f Mon Sep 17 00:00:00 2001 From: vlad-perevezentsev Date: Wed, 16 Sep 2026 17:00:53 +0200 Subject: [PATCH 8/9] Avoid a C-order copy in `dpnp.einsum` (#3069) This PR improves `dpnp.einsum` performance by avoiding a copy into C-order. Commit 039b7f14 (#3058) remaps `order="K"` to `"C"` which turned the final `dpnp.asarray` from a no-op into a full device copy because dpnp contracts operands in reverse order so `"ij,jk"` computes `(a @ b).T` and the result is naturally `f-contiguous` This PR builds the result directly in the requested layout instead of fixing it afterwards. The last contraction determines the output axis order, so it can choose which operand is on the left side of the matmul and avoid an extra copy. This is only applied for `"C"` order, since for other layouts the default operand order already matches the expected `order="K" + optimize=True` behavior. The swap does not move any data. It only changes the operand order in the matmul --- CHANGELOG.md | 1 + dpnp/dpnp_utils/dpnp_utils_einsum.py | 45 +++++++++++++++++++++++----- 2 files changed, 39 insertions(+), 7 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 0f47c903d54..1b1f37d7629 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -48,6 +48,7 @@ This release is compatible with NumPy 2.5. * `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 diff --git a/dpnp/dpnp_utils/dpnp_utils_einsum.py b/dpnp/dpnp_utils/dpnp_utils_einsum.py index efa669126a1..1f2521ef823 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] @@ -1211,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) @@ -1224,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) @@ -1245,10 +1280,6 @@ def dpnp_einsum( # a view is returned for any `order`, the same way NumPy does if not returns_view: - 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" 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) From 7808a2f7395aa6d93268228f64777f483ea5ad12 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 16 Sep 2026 10:04:26 -0700 Subject: [PATCH 9/9] Remove format strings Replaced with modern f-strings --- doc/comparison_generator.py | 44 +++++++++++----------- dpnp/dpnp_flatiter.py | 2 +- dpnp/dpnp_utils/dpnp_algo_utils.pyx | 21 +++-------- dpnp/random/dpnp_iface_random.py | 5 +-- dpnp/random/dpnp_random_state.py | 2 +- dpnp/tensor/_accumulation.py | 2 +- dpnp/tensor/_compute_follows_data.pyx | 4 +- dpnp/tensor/_dlpack.pyx | 9 ++--- dpnp/tensor/_flags.pyx | 2 +- dpnp/tensor/_indexing_functions.py | 40 ++++++-------------- dpnp/tensor/_linear_algebra_functions.py | 2 +- dpnp/tensor/_print.py | 9 ++--- dpnp/tensor/_slicing.pyx | 16 ++++---- dpnp/tests/tensor/test_usm_ndarray_capi.py | 4 +- dpnp/tests/tensor/test_usm_ndarray_ctor.py | 6 +-- 15 files changed, 66 insertions(+), 102 deletions(-) diff --git a/doc/comparison_generator.py b/doc/comparison_generator.py index 35190bec2f2..acb0dea2e9c 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_flatiter.py b/dpnp/dpnp_flatiter.py index 7375e03d802..bc17cdd22ae 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_utils/dpnp_algo_utils.pyx b/dpnp/dpnp_utils/dpnp_algo_utils.pyx index d13de1c5cdb..5a3fdd39750 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/random/dpnp_iface_random.py b/dpnp/random/dpnp_iface_random.py index 5db78086428..ca1e3d1a463 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 a18148648d7..d762c299c1f 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 069eb870f78..4def71efc58 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 798c158c7e1..65523a48c2c 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 262970385ad..9e83f46de53 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 07cae3f3e93..48adf80116c 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 48849ff7bbd..b61edb892a0 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 dcaf99b4423..c349784fdd8 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 708efc8e2ad..715a0969d2b 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 10676e63397..1d0f4afdd28 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/tests/tensor/test_usm_ndarray_capi.py b/dpnp/tests/tensor/test_usm_ndarray_capi.py index d0e44f6712d..479c288698a 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 95dd6469d1c..4e2f46d6493 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():