From aadb657e60961a2726333fbc3c727ba2e7daa06e Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Fri, 18 Sep 2026 22:57:39 -0700 Subject: [PATCH] Fix event dependencies and initial displacements The strided kernel of `full` was launched waiting only on the event which copies the packed shape and strides to the device, dropping dependent events passed to the binding removed unnecessary events added to `copy_for_reshape` `simplify_iteration_three_strides` and `simplify_iteration_four_strides` zeroed only `disp1` and `disp2`, which is not the case for other stride simplifications --- CHANGELOG.md | 3 +++ dpnp/tensor/libtensor/include/utils/strided_iters.hpp | 3 +++ dpnp/tensor/libtensor/source/copy_for_reshape.cpp | 5 +++-- dpnp/tensor/libtensor/source/full_ctor.cpp | 8 +++++++- 4 files changed, 16 insertions(+), 3 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 1b1f37d76298..eaaf2d9bc83c 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -100,6 +100,9 @@ This release is compatible with NumPy 2.5. * 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) +* Fixed the strided kernel of `dpnp.full` and `dpnp.tensor.full` not waiting on the events passed to the binding [#3073](https://github.com/IntelPython/dpnp/pull/3073) +* Fixed the list of events the copy kernel of `dpnp.reshape` and `dpnp.tensor.reshape` waits on being padded with default-constructed events [#3073](https://github.com/IntelPython/dpnp/pull/3073) +* Fixed `simplify_iteration_three_strides` and `simplify_iteration_four_strides` accumulating into their third and fourth output displacements without zeroing them first, which required the caller to initialize them [#3073](https://github.com/IntelPython/dpnp/pull/3073) ### Security diff --git a/dpnp/tensor/libtensor/include/utils/strided_iters.hpp b/dpnp/tensor/libtensor/include/utils/strided_iters.hpp index 29c17467cd78..2b3c3d706974 100644 --- a/dpnp/tensor/libtensor/include/utils/strided_iters.hpp +++ b/dpnp/tensor/libtensor/include/utils/strided_iters.hpp @@ -606,6 +606,7 @@ int simplify_iteration_three_strides(const int nd, { disp1 = StridesTy(0); disp2 = StridesTy(0); + disp3 = StridesTy(0); if (nd < 2) return nd; @@ -770,6 +771,8 @@ int simplify_iteration_four_strides(const int nd, { disp1 = StridesTy(0); disp2 = StridesTy(0); + disp3 = StridesTy(0); + disp4 = StridesTy(0); if (nd < 2) return nd; diff --git a/dpnp/tensor/libtensor/source/copy_for_reshape.cpp b/dpnp/tensor/libtensor/source/copy_for_reshape.cpp index f7c60d8cfa08..688a5763c551 100644 --- a/dpnp/tensor/libtensor/source/copy_for_reshape.cpp +++ b/dpnp/tensor/libtensor/source/copy_for_reshape.cpp @@ -152,9 +152,10 @@ std::pair const char *src_data = src.get_data(); char *dst_data = dst.get_data(); - std::vector all_deps(depends.size() + 1); - all_deps.push_back(copy_shape_ev); + std::vector all_deps; + all_deps.reserve(depends.size() + 1); all_deps.insert(std::end(all_deps), std::begin(depends), std::end(depends)); + all_deps.push_back(copy_shape_ev); sycl::event copy_for_reshape_event = fn(exec_q, src_nelems, src_nd, dst_nd, shape_strides, src_data, diff --git a/dpnp/tensor/libtensor/source/full_ctor.cpp b/dpnp/tensor/libtensor/source/full_ctor.cpp index 8345014f29b4..b2380057c5f0 100644 --- a/dpnp/tensor/libtensor/source/full_ctor.cpp +++ b/dpnp/tensor/libtensor/source/full_ctor.cpp @@ -278,9 +278,15 @@ std::pair const sycl::event ©_shape_ev = std::get<2>(ptr_size_event_tuple); py::ssize_t *shape_strides = shape_strides_owner.get(); + std::vector all_deps; + all_deps.reserve(depends.size() + 1); + all_deps.insert(std::end(all_deps), std::begin(depends), + std::end(depends)); + all_deps.push_back(copy_shape_ev); + const sycl::event &full_strided_ev = fn(exec_q, nd, dst_nelems, shape_strides, py_value, dst_data, - {copy_shape_ev}); + all_deps); // free shape_strides const auto &temporaries_cleanup_ev =