[SYCL][Matrix] Add initial get_coord API. #7037

arnamoy10 · 2022-10-12T20:13:37Z

This patch adds initial API for retrieval of coordinates from a work item element.

dkhaldi · 2022-10-13T16:02:08Z

sycl/test/matrix/matrix-bfloat16-test-coord.cpp

+
+           for (int i = 0; i < tCData.length(); ++i) {
+             size_t row, col;
+             std::tie(row, col) = tCData[i].get_coord();


you can also use size_t [ row, col] =
to avoid calling tie

dkhaldi · 2022-10-13T16:03:06Z

sycl/test/matrix/matrix-bfloat16-test-coord.cpp

+           for (int i = 0; i < tCData.length(); ++i) {
+             size_t row, col;
+             std::tie(row, col) = tCData[i].get_coord();
+             res_local_row_acc[row] += tCData[i];


you need to return res_local_row_acc and use it verify_function

dkhaldi · 2022-10-13T16:03:55Z

sycl/test/matrix/matrix-bfloat16-test-coord.cpp

+  matrix_multiply_ref((int32_t *)Aref, (int32_t *)Bref, (int32_t *)D, MATRIX_M,
+                      MATRIX_N, MATRIX_K / 2);
+
+  bool res = true;


this is what I call verify_function.
matrix_multiply_ref should also calculate sum of rows and return that instead.

dkhaldi · 2022-10-13T16:04:20Z

sycl/include/sycl/ext/oneapi/matrix/matrix-jit-use.hpp

@@ -256,6 +257,18 @@ class wi_element {
  wi_element(joint_matrix<T, NumRows, NumCols, Use, Layout, Group> &Mat,
             std::size_t i)
      : M(Mat), idx(i) {}
+
+  // Functions


remove this comment

dkhaldi · 2022-10-13T16:04:45Z

sycl/include/sycl/ext/oneapi/matrix/matrix-jit-use.hpp

@@ -256,6 +257,18 @@ class wi_element {
  wi_element(joint_matrix<T, NumRows, NumCols, Use, Layout, Group> &Mat,
             std::size_t i)
      : M(Mat), idx(i) {}
+
+  // Functions
+  std::tuple<size_t, size_t> get_coord() {


do you need to add this function to the specialization of wi_element for bfloat16 type?

dkhaldi

We also need a second test that tests all the matrices use: A, B and C. The test should not contain any mad or load operation.
Just have :
joint_matrix<bfloat16, TM, TK, use::a> sub_a(sg);
joint_matrix<bfloat16, TK, TN, use::b> sub_b(sg);
joint_matrix<float, TM, TN, use::accumulator> sub_c(sg);

the joint_matrix_fill for each of them
finally: get_coord function on each of them with the row or col sum calculation or just collecting the coordinates in a vector.

Like this, we will test this function for all three usages of the joint matrix type.

sycl/include/CL/__spirv/spirv_ops.hpp

MrSidims · 2022-10-14T16:34:56Z

sycl/include/CL/__spirv/spirv_ops.hpp

+          __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
+          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
+extern SYCL_EXTERNAL std::tuple<T, T>
+__spirv_JointMatrixWorkItemElemCoord(JOINT_MATRIX_INTEL(T, R, C, L, S, U) *,


There is no such thing as std::tuple in SPIR-V. The instruction should return int2 and if we want to create a tuple for get_coord API, then we should read elements from this vector to create tuple.

Thanks for the comment. Can you please tell a bit more about this int2 type? Is there any documentation/ code that I can take a look?

It's a 2 elements vector. int2 is a spelling from OpenCL, but guess the appropriate alias should be known for DPCPP, see: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_aliases

eh, why we use __ocl_vec_t<int32_t, 2> instead of sycl::vec<int32_t, 2> here? @MrSidims

…the vec to get the coordinates.

arnamoy10 · 2022-10-18T13:09:30Z

We also need a second test that tests all the matrices use: A, B and C. The test should not contain any mad or load operation. Just have : joint_matrix<bfloat16, TM, TK, use::a> sub_a(sg); joint_matrix<bfloat16, TK, TN, use::b> sub_b(sg); joint_matrix<float, TM, TN, use::accumulator> sub_c(sg);

the joint_matrix_fill for each of them finally: get_coord function on each of them with the row or col sum calculation or just collecting the coordinates in a vector.

Like this, we will test this function for all three usages of the joint matrix type.

Added test case

dkhaldi · 2022-10-18T13:29:10Z

sycl/include/sycl/ext/oneapi/matrix/matrix-jit-use.hpp

+
+  std::tuple<size_t, size_t> get_coord() {
+#ifdef __SYCL_DEVICE_ONLY__
+    __ocl_vec_t<int32_t, 2> co_ord =


remove underscore from the name (co_ord)

Don't see it's applied

dkhaldi · 2022-10-18T19:17:48Z

sycl/test/matrix/matrix-bfloat16-test-coord-basic.cpp

+  sycl::buffer<bfloat16, 2> bufB(B.get_data(), sycl::range<2>(K, N));
+  sycl::buffer<float, 2> bufC((float *)C.get_data(), sycl::range<2>(M, N));
+
+  sycl::buffer<int32_t, 1> res_local_row_bufA(res_local_rowA,


use usm instead of accessors

dkhaldi · 2022-10-18T19:25:24Z

sycl/test/matrix/matrix-bfloat16-test-coord-basic.cpp

+void matrix_coord_ref(int *A_mem, int *B_mem, int *C_mem, int M, int N, int K) {
+  for (int m = 0; m < M; m++)
+    for (int k = 0; k < K; k++) {
+      short *va = (short *)(A_mem + m * K + k);


change this to bfloat16 *A_mem
A_mem[m][k]

dkhaldi · 2022-10-18T19:33:04Z

sycl/test/matrix/matrix-bfloat16-test-coord-gemm.cpp

+                               N * 2, layout::packed_b);
+             sub_c = joint_matrix_mad(sg, sub_a, sub_b, sub_c);
+           }
+           joint_matrix_store(sg, sub_c,


remove the store

…at16, fix theCPU kernel

sycl/test/matrix/matrix-bfloat16-test-coord-basic.cpp

dkhaldi · 2022-10-25T13:11:52Z

sycl/test/matrix/matrix-bfloat16-test-coord-gemm.cpp

+
+           sycl::ext::oneapi::sub_group sg = spmd_item.get_sub_group();
+           joint_matrix<bfloat16, TM, TK, use::a> sub_a(sg);
+           // For B, since current implementation does not support non-packed


this comment does not apply anymore, remove it

dkhaldi · 2022-10-25T13:12:38Z

@yubingex007-a11y @MrSidims , can you please add your reviews as well?

MrSidims

@dkhaldi when do we plan to change Matrix feature macro?

MrSidims · 2022-10-25T15:05:33Z

sycl/include/CL/__spirv/spirv_ops.hpp

+          __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
+          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
+extern SYCL_EXTERNAL __ocl_vec_t<int32_t, 2>
+__spirv_JointMatrixWorkItemElemCoord(JOINT_MATRIX_INTEL(T, R, C, L, S, U) *,


Please wait with the merge until the name for the instruction is picked. In the draft SPIR-V spec version it is JointMatrixGetElementCoordINTEL

MrSidims · 2022-10-25T15:08:58Z

sycl/include/CL/__spirv/spirv_ops.hpp

+          __spv::MatrixUse U = __spv::MatrixUse::Unnecessary,
+          __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
+          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
+extern SYCL_EXTERNAL __ocl_vec_t<int32_t, 2>


For some reasons can not add suggestion.
here and after int32_t -> uint32_t

MrSidims · 2022-10-25T15:10:52Z

sycl/include/sycl/ext/oneapi/matrix/matrix-jit-use.hpp

@@ -256,6 +257,21 @@ class wi_element {
  wi_element(joint_matrix<T, NumRows, NumCols, Use, Layout, Group> &Mat,
             std::size_t i)
      : M(Mat), idx(i) {}
+
+  std::tuple<size_t, size_t> get_coord() {


nit: size_t -> uint32_t is probably better
same nit applicable to the code below

MrSidims · 2022-10-25T15:13:06Z

sycl/include/sycl/ext/oneapi/matrix/matrix-jit-use.hpp

+
+  std::tuple<size_t, size_t> get_coord() {
+#ifdef __SYCL_DEVICE_ONLY__
+    __ocl_vec_t<int32_t, 2> co_ord =


Don't see it's applied

dkhaldi · 2022-10-25T17:19:53Z

@dkhaldi when do we plan to change Matrix feature macro?

change it how?
the matrix feature macro will not change. You mean the implementation version? making use default?

dkhaldi · 2022-10-25T19:25:40Z

sycl/test/matrix/matrix-bfloat16-test-coord-basic.cpp

+             for (int i = 0; i < tAData.length(); ++i) {
+               auto [row, col] = tAData[i].get_coord();
+               resA[row] += tAData[i];
+             }


// SG size = 64, sub_a[8][16], multiple WI share a row
WI 0 --> 2 elements of row 0 --> resA[0]
WI 1 --> 3 elements of row 0 --> resA[0]
resA should be private variable --> length size
partial_sum[row] = reduction among the WI (reduce_over_group),
every WI will have same value of partial_sum[row]
copy partial_sum into the global variable

MrSidims · 2022-10-28T14:13:56Z

@dkhaldi when do we plan to change Matrix feature macro?

change it how? the matrix feature macro will not change. You mean the implementation version? making use default?

Sorry, I though that we will move feature macro with this change, but it's not the case.

sycl/include/sycl/ext/oneapi/matrix/matrix-jit-use.hpp

…pipeline is supported.

This patch adds an initial API for the retrieval of coordinates from a work item element. A `get_coord()` method is added to the intel namespace to work on `wi_element` class. Also, a relevant SPIRV op is added, which the get_coord() gets lowered to. This is recreated PR from my forked repo. The discussions are in the original (closed) PR #7037

arnamoy10 requested a review from dkhaldi October 12, 2022 20:13

arnamoy10 requested a review from a team as a code owner October 12, 2022 20:13

arnamoy10 requested a review from KseniyaTikhomirova October 12, 2022 20:13

[SYCL][Matrix] Add initial get_coord API.

59d0e7e

This patch adds initial API for retrieval of coordinates from a work item element.

arnamoy10 force-pushed the getcoordinate_support branch from 7936fb0 to 59d0e7e Compare October 13, 2022 15:09

dkhaldi requested changes Oct 13, 2022

View reviewed changes

arnamoy.bhattacharyya added 2 commits October 14, 2022 10:39

Reviewers comments

135d82b

clang-format

95adb66

dkhaldi requested changes Oct 14, 2022

View reviewed changes

sycl/include/CL/__spirv/spirv_ops.hpp Show resolved Hide resolved

MrSidims reviewed Oct 14, 2022

View reviewed changes

arnamoy.bhattacharyya added 2 commits October 17, 2022 16:59

Using olc_vec type in the spirv operation and creating a tuple using …

c7e6000

…the vec to get the coordinates.

clang-format

7742fd9

dkhaldi reviewed Oct 18, 2022

View reviewed changes

dkhaldi requested a review from yubingex007-a11y October 18, 2022 13:30

Review comments

b4e3ef5

dkhaldi reviewed Oct 18, 2022

View reviewed changes

arnamoy.bhattacharyya added 2 commits October 24, 2022 12:48

Makeaccess through USM, also update the basic kernel with use of bflo…

57a97cf

…at16, fix theCPU kernel

clang-format

1c5ace5

KseniyaTikhomirova requested a review from dkhaldi October 25, 2022 09:20

dkhaldi reviewed Oct 25, 2022

View reviewed changes

MrSidims reviewed Oct 25, 2022

View reviewed changes

dkhaldi reviewed Oct 25, 2022

View reviewed changes

arnamoy.bhattacharyya added 2 commits November 4, 2022 11:02

Reviewer comments

4529e66

Clang-format

643aafc

arnamoy.bhattacharyya added 2 commits November 4, 2022 11:09

More comments addressed.

8e73c9d

Fixing small error

b2ca8e4

yubingex007-a11y reviewed Nov 9, 2022

View reviewed changes

sycl/include/sycl/ext/oneapi/matrix/matrix-jit-use.hpp Show resolved Hide resolved

Adding XFAIL to test cases when we run. Will take away when the full …

d2dfda6

…pipeline is supported.

whitneywhtsang closed this Dec 1, 2022

whitneywhtsang deleted the getcoordinate_support branch December 1, 2022 19:59

This was referenced Dec 13, 2022

[SYCL][Matrix] Add initial get_coord API. #7772

Closed

[SYCL][Matrix] Add initial get_coord API #7851

Merged

[SYCL][Matrix] Add initial get_coord API. #7037

[SYCL][Matrix] Add initial get_coord API. #7037

Uh oh!

Conversation

arnamoy10 commented Oct 12, 2022

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

dkhaldi left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

yubingex007-a11y Oct 26, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

arnamoy10 commented Oct 18, 2022

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

dkhaldi commented Oct 25, 2022

Uh oh!

MrSidims left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

dkhaldi commented Oct 25, 2022

Uh oh!

yubingex007-a11y Oct 26, 2022 •

edited

Loading