Skip to content

[SYCL][Matrix spec] keep deletion of assign op and copy ctor but change signature of joint_matrix_mad #11007

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

Merged
steffenlarsen merged 2 commits into from
Sep 18, 2023
Merged

[SYCL][Matrix spec] keep deletion of assign op and copy ctor but change signature of joint_matrix_mad #11007

Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
30eeaba
[SYCL][Matrix spec] Do not delete assign op and copy ctor as they are…
dkhaldi Aug 29, 2023
8e139ed
Keep ctors deleted, make D input in mad, remove deprecated API like m…
dkhaldi Sep 15, 2023
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
13 changes: 7 additions & 6 deletions sycl/doc/extensions/experimental/sycl_ext_matrix/sycl_ext_oneapi_matrix.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -109,16 +109,17 @@ template <typename Group, typename T, use Use, size_t Rows, size_t Cols,
layout::dynamic : /*unspecified*/ >
struct joint_matrix {
joint_matrix();
joint_matrix(const joint_matrix &) = delete;
joint_matrix &operator=(const joint_matrix &) = delete;
joint_matrix(const joint_matrix &);
joint_matrix &operator=(const joint_matrix &);
};

} // namespace sycl::ext::oneapi::experimental::matrix
```
The constructor for the `joint_matrix` type is a group function as
defined in section 4.17.3 of the core SYCL specification. It must be
encountered in converged control flow by all work-items in the
`Group`.
The constructors for the `joint_matrix` type and the assignment
operator are group functions as defined in section 4.17.3 of the core
SYCL specification. They must be encountered in converged control flow
by all work-items in the `Group`. Note that the assignment operator
and the copy constructor do not copy the entire matrix content.
Copy link
Contributor

@gmlueck gmlueck Sep 6, 2023

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why do we need to say that the copy constructor and assignment operator don't copy the matrix content? How would the application even be able to tell whether the matrix content was copied? Isn't the content of the matrix hidden behind this abstract joint_matrix type?

I'm thinking now that we could replace joint_matrix_copy with an assignment operator that allows the user to copy from a matrix with different Use and T (type). That way, matrix copy operations will always be done via the assignment operator. That seems clearer than having both an assignment operator and a joint_matrix_copy function.

Copy link
Contributor Author

@dkhaldi dkhaldi Sep 8, 2023 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@gmlueck, if we have the assignment op do the actual data copy, = joint_matrix_mad that uses assignment op would do extra unnecessary copy of the data (

llvm/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp

Line 418 in faec54c

joint_matrix_mad(
)

Copy link
Contributor

@gmlueck gmlueck Sep 11, 2023

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It seems like neither joint_matrix_copy nor the assignment operator should copy anything to/from memory. The only memory operations are joint_matrix_load and joint_matrix_store. It seems like both joint_matrix_copy and the assignment operator should both just copy some "handle" that represents the register(s) that hold the matrix content?

In your implementation today, what is the difference between the implementation of joint_matrix_copy and the copy operator?

Here's another way to think about it. Consider a code snippet like this:

joint_matrix<sub_group, int8_t, use::a, Rows, Cols, layout::row_major> m1;
joint_matrix<sub_group, int8_t, use::a, Rows, Cols, layout::row_major> m2;
joint_matrix_load(sg, m1, ptr1, Rows);  // Load "m1" with data from "ptr1"
m2 = m1;
joint_matrix_load(sg, m1, ptr2, Rows);  // Load "m1" with data from "ptr2"

I think a user would expect the matrix m2 to still contain the data loaded from ptr1 after the end of that code snippet. Isn't this the same semantic as joint_matrix_copy?


==== Group Memory Scope
Most operations on the joint_matrix are group functions, meaning that
Expand Down