From a6c19b41e16c098974f47d4e0eb2bbbfd421619a Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 12 Aug 2026 14:37:29 +0200 Subject: [PATCH 1/2] Guard topk against k == 0 top_k(x, 0) is a valid request whose result is empty, but py_topk still dispatched to the kernel, which builds an nd_range with a zero global size (nelems = iter_nelems * k == 0) and a non-zero local size. That is a silent no-op on runtimes built with NDEBUG, but aborts on an assertions-enabled SYCL runtime (adjustNDRangePerKernel asserts NDR.LocalSize[0] == 0 when GlobalSize is 0). Extend the early-return guard to cover k == 0. The vals/inds outputs are already allocated with a zero-length result axis, so there is nothing to compute or fill. --- dpnp/tensor/libtensor/source/sorting/topk.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dpnp/tensor/libtensor/source/sorting/topk.cpp b/dpnp/tensor/libtensor/source/sorting/topk.cpp index 0103839f3004..db4b99ab879b 100644 --- a/dpnp/tensor/libtensor/source/sorting/topk.cpp +++ b/dpnp/tensor/libtensor/source/sorting/topk.cpp @@ -219,7 +219,7 @@ std::pair dpnp::tensor::validation::CheckWritable::throw_if_not_writable(vals); dpnp::tensor::validation::CheckWritable::throw_if_not_writable(inds); - if ((iter_nelems == 0) || (axis_nelems == 0)) { + if ((iter_nelems == 0) || (axis_nelems == 0) || (k == 0)) { // Nothing to do return std::make_pair(sycl::event(), sycl::event()); } From 588e039ae1a1ff3be865c57db5b4a23b0a2fa31e Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 12 Aug 2026 14:38:06 +0200 Subject: [PATCH 2/2] Update changelog for topk k == 0 fix --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 8d01978e60df..7dee9a8a90e9 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -78,6 +78,7 @@ This release is compatible with NumPy 2.5. * Fixed `dpnp.interp` returning `nan` when querying at an exact knot point whose adjacent `fp` value is `inf` [#2986](https://github.com/IntelPython/dpnp/pull/2986) * Fixed missing strides validation in `dpnp.tensor.usm_ndarray` constructor when allocating new memory [#2927](https://github.com/IntelPython/dpnp/pull/2927) * Fixed `dpnp.bincount` raising a `ValueError` on an empty input array instead of returning an empty `intp` array [#3018](https://github.com/IntelPython/dpnp/pull/3018) +* Fixed `dpnp.tensor.top_k` aborting for `k=0` by returning empty result arrays without launching a zero-sized kernel [#3022](https://github.com/IntelPython/dpnp/pull/3022) ### Security