From 414422bec3ff151059f4297183fafbbdccd5ee2d Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Tue, 8 Sep 2026 14:06:48 +0200 Subject: [PATCH 1/4] Fix row corruption in batched multi-level cumulative scan The intermediate block-scan update in inclusive_scan_iter used an unstrided (NoOp) row offset into the strided per-row scan buffers, so cumulative_sum/cumulative_prod along an axis returned wrong results whenever the scan needed >=3 reduction levels (axis > chunk_size**2) with more than one batch row. Use a Strided1DIndexer with the per-row stride instead. Add a regression test exercising the >=3-level batched scan path. --- .../libtensor/include/kernels/accumulators.hpp | 6 ++++-- dpnp/tests/tensor/test_tensor_accumulation.py | 17 +++++++++++++++++ 2 files changed, 21 insertions(+), 2 deletions(-) diff --git a/dpnp/tensor/libtensor/include/kernels/accumulators.hpp b/dpnp/tensor/libtensor/include/kernels/accumulators.hpp index 079fcf5e9c96..fe3cec59a179 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 b7ea9147e100..57863e8acd07 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): + # axis > chunk_size**2 (chunk_size <= 2048) needs >=3 scan levels; with >1 + # batch row the intermediate update mis-strided rows. 5e6 forces that path. + 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 50ce2e5df6544c4a6031522427ed7e774430cf6c Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Tue, 8 Sep 2026 14:13:05 +0200 Subject: [PATCH 2/4] Add changelog entry for the cumulative scan fix --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 09961c7ccd34..6216422546c5 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -95,6 +95,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 a axis of an array with more than one row [#3063](https://github.com/IntelPython/dpnp/pull/3063) ### Security From eb084685ae1c2c6e7fb5d067f2e4b5049ee59b31 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Tue, 8 Sep 2026 14:53:32 +0200 Subject: [PATCH 3/4] Fix grammar in cumulative scan changelog entry --- CHANGELOG.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 6216422546c5..bf86dce28f4e 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -95,7 +95,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 a axis of an array with more than one row [#3063](https://github.com/IntelPython/dpnp/pull/3063) +* 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 From a1376ed2389bb86b2a52b94802b9d5aa79251349 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Tue, 8 Sep 2026 17:33:02 +0200 Subject: [PATCH 4/4] Note batched multi-level scan test as a regression test --- dpnp/tests/tensor/test_tensor_accumulation.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/dpnp/tests/tensor/test_tensor_accumulation.py b/dpnp/tests/tensor/test_tensor_accumulation.py index 57863e8acd07..d37cb706e3e5 100644 --- a/dpnp/tests/tensor/test_tensor_accumulation.py +++ b/dpnp/tests/tensor/test_tensor_accumulation.py @@ -128,8 +128,8 @@ def test_strided_cumsum_axis_sint(dt): @pytest.mark.parametrize("func", ["cumulative_sum", "cumulative_prod"]) def test_batched_multilevel_scan(func): - # axis > chunk_size**2 (chunk_size <= 2048) needs >=3 scan levels; with >1 - # batch row the intermediate update mis-strided rows. 5e6 forces that path. + # 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")