Fix row corruption in batched multi-level cumulative scan

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
 

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<UpdateKernelName, outputT, n_wi,
-                                     NoOpIndexer, NoOpIndexer, ScanOpT>(
+                                     Strided1DIndexer, NoOpIndexer, ScanOpT>(
         exec_q, iter_nelems, src, src_size, local_scans, chunk_size,
         local_stride, iter_out_indexer, out_indexer, dependent_event);
 }

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()