diff --git a/CHANGELOG.md b/CHANGELOG.md index e1715ec8327..7136ebc7e53 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -101,6 +101,9 @@ This release is compatible with NumPy 2.5. * 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 `dpnp.ndarray` constructor returning a view at the wrong address [#3068](https://github.com/IntelPython/dpnp/pull/3068) +* 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 kernels of `dpnp.reshape`, `dpnp.tensor.reshape`, `dpnp.roll` and `dpnp.tensor.roll` wait 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 29c17467cd7..2b3c3d70697 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 f7c60d8cfa0..688a5763c55 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/copy_for_roll.cpp b/dpnp/tensor/libtensor/source/copy_for_roll.cpp index 662a0390e0b..4ff6cca14e4 100644 --- a/dpnp/tensor/libtensor/source/copy_for_roll.cpp +++ b/dpnp/tensor/libtensor/source/copy_for_roll.cpp @@ -233,9 +233,10 @@ std::pair sycl::event copy_shape_ev = std::get<2>(ptr_size_event_tuple); const py::ssize_t *shape_strides = shape_strides_owner.get(); - 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_roll_event = fn(exec_q, offset, src_nelems, src_nd, shape_strides, src_data, @@ -357,9 +358,10 @@ std::pair sycl::event copy_shape_ev = std::get<2>(ptr_size_event_tuple); const py::ssize_t *shape_strides_shifts = shape_strides_shifts_owner.get(); - 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_roll_event = fn(exec_q, src_nelems, src_nd, shape_strides_shifts, src_data, diff --git a/dpnp/tensor/libtensor/source/full_ctor.cpp b/dpnp/tensor/libtensor/source/full_ctor.cpp index 8345014f29b..b2380057c5f 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 =