Enable matrix_load, matrix_store, and matrix_mad #4076

yubingex007-a11y · 2021-07-08T12:52:22Z

Change-Id: I97b939c5559317ad82bec460221bccb83975b004

bader

Please, put files to ext/intel directory instead of ONEAPI.

dkhaldi · 2021-07-19T14:31:47Z

Please, put files to ext/intel directory instead of ONEAPI.

@bader, Although the matrix feature is in experimental phase, it is a oneAPI feature, it is not an intel specific feature. I think we should keep it as it was already.

bader · 2021-07-19T14:37:23Z

Sorry, I meant to say ext/oneapi. Please, rebase on the tip of the branch to resolve the conflicts.
See d703f57 for more details on the requested change. We moved all oneAPI extensions from ONEAPI to ext/oneapi directory.

dkhaldi · 2021-07-19T14:55:02Z

Sorry, I meant to say ext/oneapi. Please, rebase on the tip of the branch to resolve the conflicts.
See d703f57 for more details on the requested change. We moved all oneAPI extensions from ONEAPI to ext/oneapi directory.

What would be the namespace we should use in this case?
Currently we are using: sycl:ext::intel::experimental::matrix;
Should it be: sycl:ext::ONEAPI::experimental::matrix;?
BTW, I see experimental directory in ONEAPI/. Should we move matrix/ there?

bader · 2021-07-19T15:07:34Z

The namespace should be sycl::ext::oneapi::experimental::matrix.

According to my understanding ONEAPI directory is left for backward compatibility only and should not be used for new features.
All new features should follow SYCL spec guidelines and use sycl::ext::oneapi namespace (implementation guideline is to put files into ext/oneapi directory reflecting namespace structure).

Tagging @dm-vodopyanov and @romanovvlad to correct me if I miss something.

romanovvlad · 2021-07-20T08:05:34Z

The namespace should be sycl::ext::oneapi::experimental::matrix.

According to my understanding ONEAPI directory is left for backward compatibility only and should not be used for new features.
All new features should follow SYCL spec guidelines and use sycl::ext::oneapi namespace (implementation guideline is to put files into ext/oneapi directory reflecting namespace structure).

Tagging @dm-vodopyanov and @romanovvlad to correct me if I miss something.

Please, use sycl::ext::oneapi::experimental for general extension and sycl::ext::intel::experimental if the extension is specific to Intel HW.

yubingex007-a11y · 2021-07-22T07:14:19Z

sycl/include/sycl/ext/oneapi/matrix/matrix-aot-amx.hpp

@@ -38,7 +38,7 @@
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace ext {
-namespace intel {
+namespace oneapi {


BTW, @bader @dkhaldi , should we move this file to ext/intel directory since it works on intel HW?

mistakenly change it, sorry for inconvience.

BTW, @bader @dkhaldi , should we move this file to ext/intel directory since it works on intel HW?

Please keep everything in ext/oneapi.
The feature is still oneAPI "matrix". It is just that the initial implementation targeted only AMX.

1. change namespace to ext::intel 2. move matrix-aot-amx.hpp to ext/intel directory

AlexeySotkin · 2021-07-22T13:59:43Z

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

+          matrix_layout Layout = matrix_layout::row_major,
+          access::address_space Space>
+inline __SYCL_ALWAYS_INLINE void joint_matrix_store(
+    Group sg, joint_matrix<Group, T, NumRows, NumCols, Layout> &res,


Suggested change

Group sg, joint_matrix<Group, T, NumRows, NumCols, Layout> &res,

Group sg, joint_matrix<Group, T, NumRows, NumCols, Layout> &obj,

Do we need to change joint_matrix_load as well?

@AlexeySotkin can you please elaborate why you asked for the change res --> obj?

Just improving the name. There is no result for the store operation. It returns void. This parameter is a matrix object which we want to store.

AlexeySotkin · 2021-07-22T14:22:29Z

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

+          access::address_space Space>
+inline __SYCL_ALWAYS_INLINE void joint_matrix_load(
+    Group sg, joint_matrix<Group, T, NumRows, NumCols, Layout> &res,
+    multi_ptr<T, Space> src, size_t stride, matrix_layout L = Layout) {


Function argument L is not used. Probably we should use it as operand of __spirv_MatrixLoadINTEL, but it must be a constant.
It seems like we have two layouts here: one for the memory we read from, and another for the matrix object, right? If so, I think we need another template parameter to specify layout of the memory.

These are the changes we agreed on in our meeting today:

remove matrix_layout L = Layout

pass L (the memory layout) to the load/store functions

rename L --> memL

rename Layout --> matL

rename the header file spirv API from matrix to joint_matrix

It is not quite clear what is "pass L to the load/store functions". I think we should move the last function parameter to template parameters, right?
Like:

template <typename Group, typename T, size_t NumRows, size_t NumCols, matrix_layout MemL = matrix_layout::row_major, matrix_layout MatL = matrix_layout::row_major, access::address_space Space> inline __SYCL_ALWAYS_INLINE void joint_matrix_load( Group sg, joint_matrix<Group, T, NumRows, NumCols, Layout> &res, multi_ptr<T, Space> src, size_t stride) { ... res.spvm = __spirv_MatrixLoadINTEL<T, NumRows, NumCols, spv_matrix_layout_traits<MemL>::value, spv_matrix_layout_traits<MatL>::value>(Ptr, stride); ... }

Right now, the user does not need to specify any template argument when calling load/store. If we move the memory layout to the template, the user will have to specify each of the parameters (deduce all or nothing) when calling load/store.

The other reason is that I would like to not converge too much (unless there is a big reason) from Vulkan and WMMA APIs thus keeping the meory layout as argument to the functions

yubingex007-a11y · 2021-07-23T13:36:28Z

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

+  T *Ptr = src.get();
+  res.spvm = __spirv_MatrixLoadINTEL<T, NumRows, NumCols,
+                                     spv_matrix_layout_traits<Layout>::value>(
+      Ptr, stride, spv_matrix_layout_traits<Layout>::value);


yubingex007-a11y · 2021-07-23T13:37:20Z

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

+};
+
+template <typename Group, typename T, size_t NumRows, size_t NumCols,
+          matrix_layout Layout = matrix_layout::row_major,


dkhaldi · 2021-07-26T14:52:30Z

sycl/test/matrix/matrix-amx-bf16-test.cpp

@@ -1,4 +1,4 @@
-// RUN: %clangxx -march=sapphirerapids -fsycl -O2 %s -o %t.out
+// RUN: %clangxx -DSYCL_EXT_ONEAPI_MATRIX=1 -march=sapphirerapids -fsycl -O2 %s -o %t.out
 #include <CL/sycl.hpp>
 #if (SYCL_EXT_ONEAPI_MATRIX == 1)


// Only the leader perform AMX computation.
Why do you have "the leader" test in the code. this is wrong for two reasons:
1- SG size is one so we don't need to test that.
2- all work items in the subgroup must enter the joint matrix code. So we should NOT have diverged code in there

dkhaldi · 2021-07-26T14:54:20Z

sycl/test/matrix/matrix-bf16-test-SG-16.cpp

+
+           // Only the leader perform AMX computation.
+           if (spmd_item.get_local_id(1) % TILE_SZ)
+             return;


Again, we should not have "the leader" test in the code. this is wrong because all work items in the subgroup must enter the joint matrix code. So we should not have diverge code in there.
The OpenCL backend will not vectorize the code so we should be able to generate ONE AMX code for the whole sub group

dkhaldi · 2021-07-26T14:55:09Z

sycl/test/matrix/matrix-bf16-test.cpp

+
+           // Only the leader perform AMX computation.
+           if (spmd_item.get_local_id(1) % TILE_SZ)
+             return;


same comment as above

dkhaldi · 2021-07-26T14:56:24Z

sycl/test/matrix/matrix-int8-test-SG-16.cpp

+
+           // Only the leader perform AMX computation.
+           if (spmd_item.get_local_id(1) % TILE_SZ)
+             return;


same comment as above

dkhaldi · 2021-07-26T14:57:17Z

sycl/test/matrix/matrix-int8-test.cpp

+
+           // Only the leader perform AMX computation.
+           if (spmd_item.get_local_id(1) % TILE_SZ)
+             return;


same comment as above

dkhaldi · 2021-08-03T19:40:55Z

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

+  case matrix_layout::row_major:
+    res.spvm =
+        __spirv_JointMatrixLoadINTEL<T, NumRows, NumCols,
+                                     spv_matrix_layout_traits<Layout>::value>(


you omitted the group argument here right?

dkhaldi · 2021-08-03T19:41:44Z

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

+    res.spvm =
+        __spirv_JointMatrixLoadINTEL<T, NumRows, NumCols,
+                                     spv_matrix_layout_traits<Layout>::value>(
+            Ptr, stride, __spv::MatrixLayout::ColumnMajor);


same as above

dkhaldi · 2021-08-03T19:42:48Z

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

+    res.spvm =
+        __spirv_JointMatrixLoadINTEL<T, NumRows, NumCols,
+                                     spv_matrix_layout_traits<Layout>::value>(
+            Ptr, stride, __spv::MatrixLayout::PackedA);


spirv should have scope argument so sg should be passed in all load/store/mad functions

dkhaldi · 2021-08-03T19:43:35Z

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

+  case matrix_layout::col_major:
+    __spirv_JointMatrixStoreINTEL<T, NumRows, NumCols,
+                                  spv_matrix_layout_traits<MatL>::value>(
+        Ptr, obj.spvm, stride, __spv::MatrixLayout::ColumnMajor);


I thought we said obj should be change to "src" right?

dkhaldi · 2021-08-10T14:44:18Z

LGTM

My comment is addressed. Thanks.

bader · 2021-08-10T17:24:02Z

LGTM

@dkhaldi, you must approve using GitHub UI to unblock the merge. Currently it's blocked by your previous review where changes were requested.

dkhaldi · 2021-08-10T17:51:35Z

LGTM

@dkhaldi, you must approve using GitHub UI to unblock the merge. Currently it's blocked by your previous review where changes were requested.

@bader done

yubingex007-a11y · 2021-08-11T14:47:54Z

@intel/llvm-reviewers-runtime Hi, Could you do a code review please?

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

1- we move the Group template argument on joint_matrix to the last position and make it optional. By default set Group to subgroup. So the user does not need to specify this each time 2- we change the "mad" function name from joint_matrix_mad to joint_matrix_mma.

dkhaldi · 2021-08-13T15:07:17Z

LGTM, @bader can you help merge this PR?

bader · 2021-08-13T16:14:28Z

LGTM, @bader can you help merge this PR?

The PR is waiting on code owner review from @intel/llvm-reviewers-runtime.

s-kanaev

matrix-amx.hpp => matrix-aot-amx.hpp
Does this renaming follow API/ABI compatibility policy?

sycl/include/CL/sycl/ONEAPI/matrix/matrix-amx.hpp said one should include matrix-amx.hpp and the latter file becomes unavailable now.

Probably, the user won't be happy.

Other than that, the changes seem to be legit.

yubingex007-a11y · 2021-08-16T14:38:03Z

matrix-amx.hpp => matrix-aot-amx.hpp
Does this renaming follow API/ABI compatibility policy?

sycl/include/CL/sycl/ONEAPI/matrix/matrix-amx.hpp said one should include matrix-amx.hpp and the latter file becomes unavailable now.

Probably, the user won't be happy.

Other than that, the changes seem to be legit.
@s-kanaev , Thanks for your review.
This is an experimental interface so it is assumed to be changed.

s-kanaev

LGTM

Enable matrix_load, matrix_store, and matrix_mad

9f55ca8

Change-Id: I97b939c5559317ad82bec460221bccb83975b004

yubingex007-a11y requested review from dkhaldi and LuoYuanke July 8, 2021 12:52

yubingex007-a11y requested a review from a team as a code owner July 8, 2021 12:52

yubingex007-a11y requested a review from sergey-semenov July 8, 2021 12:52

Delete some useless changes

a19b386

dkhaldi requested review from Pennycook, AlexeySotkin and bader July 14, 2021 15:58

bader previously requested changes Jul 18, 2021

View reviewed changes

Merge remote-tracking branch 'origin/sycl' into matrixv2

9fd9334

yubingex007-a11y commented Jul 22, 2021

View reviewed changes

Bugfix for matrix-aot-amx:

3229cdd

1. change namespace to ext::intel 2. move matrix-aot-amx.hpp to ext/intel directory

AlexeySotkin reviewed Jul 22, 2021

View reviewed changes

yubingex007-a11y commented Jul 23, 2021

View reviewed changes

dkhaldi reviewed Jul 26, 2021

View reviewed changes

yubingex007-a11y added 3 commits July 28, 2021 15:53

Address Dounia's comments

2fc15f5

Fix issues hinted by clang-format

acbcc75

Fix the issue about MemL

8705072

yubingex007-a11y requested a review from dkhaldi August 3, 2021 19:31

dkhaldi requested changes Aug 3, 2021

View reviewed changes

define traits for scope

b3243d9

yubingex007-a11y requested a review from dkhaldi August 5, 2021 06:35

yubingex007-a11y requested a review from bader August 10, 2021 14:44

dkhaldi approved these changes Aug 10, 2021

View reviewed changes

yubingex007-a11y requested a review from AlexeySotkin August 11, 2021 07:11

dkhaldi requested changes Aug 11, 2021

View reviewed changes

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

yubingex007-a11y added 2 commits August 12, 2021 22:28

Go back to joint_matrix_mad

8a30b1e

dkhaldi approved these changes Aug 13, 2021

View reviewed changes

s-kanaev reviewed Aug 16, 2021

View reviewed changes

s-kanaev approved these changes Aug 17, 2021

View reviewed changes

romanovvlad merged commit 7f21853 into intel:sycl Aug 17, 2021

	Group sg, joint_matrix<Group, T, NumRows, NumCols, Layout> &res,
	Group sg, joint_matrix<Group, T, NumRows, NumCols, Layout> &obj,

Enable matrix_load, matrix_store, and matrix_mad #4076

Enable matrix_load, matrix_store, and matrix_mad #4076

Uh oh!

Conversation

yubingex007-a11y commented Jul 8, 2021

Uh oh!

bader left a comment

Choose a reason for hiding this comment

Uh oh!

dkhaldi commented Jul 19, 2021

Uh oh!

bader commented Jul 19, 2021

Uh oh!

dkhaldi commented Jul 19, 2021

Uh oh!

bader commented Jul 19, 2021

Uh oh!

romanovvlad commented Jul 20, 2021

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 Jul 22, 2021 • 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!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

yubingex007-a11y Jul 23, 2021 • 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!

yubingex007-a11y Jul 23, 2021 • 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!

dkhaldi Jul 26, 2021 • 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!

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 commented Aug 10, 2021

Uh oh!

bader commented Aug 10, 2021

Uh oh!

dkhaldi commented Aug 10, 2021

Uh oh!

yubingex007-a11y commented Aug 11, 2021

Uh oh!

Uh oh!

dkhaldi Jul 22, 2021 •

edited

Loading

yubingex007-a11y Jul 23, 2021 •

edited

Loading

yubingex007-a11y Jul 23, 2021 •

edited

Loading

dkhaldi Jul 26, 2021 •

edited

Loading