Skip to content

Fix event dependency bugs and initial displacement initialization #3072

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
ndgrigorian wants to merge 2 commits into
base: master
Choose a base branch
from
Open

Fix event dependency bugs and initial displacement initialization #3072

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
3389a93
Fix event dependencies and initial displacements
ndgrigorian Sep 19, 2026
fe516a1
Fix padding the dependencies of copy_for_roll
ndgrigorian Sep 25, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 3 additions & 0 deletions CHANGELOG.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand Down
3 changes: 3 additions & 0 deletions dpnp/tensor/libtensor/include/utils/strided_iters.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;

Expand Down Expand Up @@ -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;

Expand Down
5 changes: 3 additions & 2 deletions dpnp/tensor/libtensor/source/copy_for_reshape.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -152,9 +152,10 @@ std::pair<sycl::event, sycl::event>
const char *src_data = src.get_data();
char *dst_data = dst.get_data();

std::vector<sycl::event> all_deps(depends.size() + 1);
all_deps.push_back(copy_shape_ev);
std::vector<sycl::event> all_deps;
all_deps.reserve(depends.size() + 1);
all_deps.insert(std::end(all_deps), std::begin(depends), std::end(depends));
Comment thread
ndgrigorian marked this conversation as resolved.
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,
Expand Down
10 changes: 6 additions & 4 deletions dpnp/tensor/libtensor/source/copy_for_roll.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -233,9 +233,10 @@ std::pair<sycl::event, sycl::event>
sycl::event copy_shape_ev = std::get<2>(ptr_size_event_tuple);
const py::ssize_t *shape_strides = shape_strides_owner.get();

std::vector<sycl::event> all_deps(depends.size() + 1);
all_deps.push_back(copy_shape_ev);
std::vector<sycl::event> 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,
Expand Down Expand Up @@ -357,9 +358,10 @@ std::pair<sycl::event, sycl::event>
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<sycl::event> all_deps(depends.size() + 1);
all_deps.push_back(copy_shape_ev);
std::vector<sycl::event> 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,
Expand Down
8 changes: 7 additions & 1 deletion dpnp/tensor/libtensor/source/full_ctor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -278,9 +278,15 @@ std::pair<sycl::event, sycl::event>
const sycl::event &copy_shape_ev = std::get<2>(ptr_size_event_tuple);
py::ssize_t *shape_strides = shape_strides_owner.get();

std::vector<sycl::event> 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 =
Expand Down
Loading