[SYCL][Matrix] Enable wi_slice for joint_matrix #4979

yubingex007-a11y · 2021-11-17T16:43:24Z

No description provided.

yubingex007-a11y · 2021-11-17T16:45:38Z

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

+            auto wi_slice_c = joint_matrix_get_slice(sub_c); // M.get_wi_slice()
+            for (int i = 0; i < wi_slice.length(); i++) {
+                wi_slice_c.data[i] *= 1;
+            }


I just create a testcase from previous case and here I do sub_c *1=1 with wi_slice.

dkhaldi · 2021-11-17T19:02:57Z

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

+                             N, matrix_layout::row_major);
+            auto wi_slice_c = joint_matrix_get_slice(sub_c); // M.get_wi_slice()
+            for (int i = 0; i < wi_slice.length(); i++) {
+                wi_slice_c.data[i] *= 1;


this should look like this: wi_slice_c[i] *=1;
is this because you were not able to overload directly operator[]?

We can't.
Please see http://cmplrexplorer.intel.com/z/PGc96P

dkhaldi · 2021-11-17T19:03:48Z

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

+                             accC.get_pointer() + (sg_startx * TM) * N +
+                                 sg_starty / SG_SZ * TN,
+                             N, matrix_layout::row_major);
+            auto wi_slice_c = joint_matrix_get_slice(sub_c); // M.get_wi_slice()


change the name to get_work_item_data as this is not a "joint" instruction

dkhaldi · 2021-11-24T14:37:52Z

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

+                             accC.get_pointer() + (sg_startx * TM) * N +
+                                 sg_starty / SG_SZ * TN,
+                             N, matrix_layout::row_major);
+            auto wi_slice_c = sub_c.get_wi_slice(); // M.get_wi_slice()


change the name of the function to >> get_wi_data

dkhaldi · 2021-11-24T14:38:36Z

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

+                              accC.get_pointer() + (sg_startx * TM) * N +
+                                  sg_starty / SG_SZ * TN,
+                              N, matrix_layout::row_major);
+         }); // parallel for


1st: change the position of the mul
C = mad()
for (int i = 0; i < wi_slice_c.length(); i++) {
wi_slice_c[i] *= alpha;
}
2- change the name of the test: matrix-elemwise-ops.cpp

dkhaldi · 2021-11-29T15:47:32Z

sycl/test/matrix/matrix-elemwise-ops.cpp

@@ -86,6 +85,10 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C, big_matrix<T2, N
                                   sg_starty / SG_SZ * TN * 4,
                               N * 4, matrix_layout::packed_b);
             sub_c = joint_matrix_mad(sg, sub_a, sub_b, sub_c);
+             auto wi_slice_c = sub_c.get_wi_data(); // M.get_wi_data()


remove the comment

AlexeySotkin · 2021-12-06T11:45:28Z

sycl/include/CL/__spirv/spirv_ops.hpp

+
+template <typename T, std::size_t R, std::size_t C, __spv::MatrixLayout U,
+          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
+extern SYCL_EXTERNAL size_t __spirv_JointMatrixGetSliceLength(


Should we add Scope parameter to this function, like we did for __spirv_JointMatrixMadINTEL?

yubingex007-a11y · 2021-12-14T14:07:44Z

ping?

dkhaldi · 2021-12-14T20:28:54Z

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

+  }
+  wi_elem &operator=(const T &rhs) {
+#ifdef __SYCL_DEVICE_ONLY__
+    M.spvm = __spirv_JointMatrixSetSliceElem(M.spvm, idx, rhs);


Note that these will be converted to
SPIRVVectorInsertDynamic
and
SPIRVVectorextractDynamic
in SPIRV.
and to:
int_experimental_matrix_wi_slice_extract
int_experimental_matrix_wi_slice_insert
in LLVM IR
so we need to change the names you are using to insert/extract instead of set/get. Something like:
__spirv_JointMatrix_Extract
__spirv_JointMatrix_Insert
is sufficient to match the SPIRV ones.
What do you think @AlexeySotkin?

We should change it to

Suggested change

M.spvm = __spirv_JointMatrixSetSliceElem(M.spvm, idx, rhs);

M.spvm = __spirv_VectorInsertDynamic(M.spvm, rhs, idx);

dkhaldi · 2021-12-14T20:33:59Z

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

+template <typename T, size_t NumRows, size_t NumCols,
+          matrix_layout Layout = matrix_layout::row_major,
+          typename Group = sycl::sub_group>
+class wi_elem {


change it to wi_element

AlexeySotkin · 2021-12-15T14:42:57Z

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

+    M.spvm = __spirv_JointMatrixSetSliceElem(
+        M.spvm, idx, __spirv_JointMatrixGetSliceElem(M.spvm, idx) * rhs);


Suggested change

M.spvm = __spirv_JointMatrixSetSliceElem(

M.spvm, idx, __spirv_JointMatrixGetSliceElem(M.spvm, idx) * rhs);

M.spvm = __spirv_VectorInsertDynamic(

M.spvm, __spirv_VectorExtractDynamic(M.spvm, idx) * rhs, idx,);

AlexeySotkin · 2021-12-15T14:43:57Z

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

+  wi_slice(joint_matrix<T, NumRows, NumCols, Layout, Group> &Mat) : M(Mat) {}
+  size_t length() {
+#ifdef __SYCL_DEVICE_ONLY__
+    return __spirv_JointMatrixGetSliceLength(M.spvm);


Suggested change

return __spirv_JointMatrixGetSliceLength(M.spvm);

return __spirv_JointMatrixWorkItemLengthINTEL(M.spvm);

AlexeySotkin · 2021-12-15T14:45:31Z

sycl/include/CL/__spirv/spirv_ops.hpp

+template <typename T>
+using __spirv_wi_slice_t = T __attribute__((ext_vector_type(0xffffff)));
+
+template <typename T, std::size_t R, std::size_t C, __spv::MatrixLayout U,
+          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
+extern SYCL_EXTERNAL __spirv_wi_slice_t<T> &__spirv_JointMatrixGetSliceData(
+    __spv::__spirv_JointMatrixINTEL<T, R, C, U, S> *);


With the current SPIR-V design we don't need these.

AlexeySotkin · 2021-12-15T14:47:17Z

sycl/include/CL/__spirv/spirv_ops.hpp

+
+template <typename T, std::size_t R, std::size_t C, __spv::MatrixLayout U,
+          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
+extern SYCL_EXTERNAL T __spirv_JointMatrixGetSliceElem(


Suggested change

extern SYCL_EXTERNAL T __spirv_JointMatrixGetSliceElem(

extern SYCL_EXTERNAL T __spirv_VectorExtractDynamic(

AlexeySotkin · 2021-12-15T14:48:59Z

sycl/include/CL/__spirv/spirv_ops.hpp

+__spirv_JointMatrixSetSliceElem(
+    __spv::__spirv_JointMatrixINTEL<T, R, C, U, S> *, size_t i, T val);


Suggested change

__spirv_JointMatrixSetSliceElem(

__spv::__spirv_JointMatrixINTEL<T, R, C, U, S> *, size_t i, T val);

__spirv_VectorInsertDynamic(

__spv::__spirv_JointMatrixINTEL<T, R, C, U, S> *, T val, size_t i);

AlexeySotkin · 2021-12-15T14:53:18Z

sycl/include/CL/__spirv/spirv_ops.hpp

+template <typename T, std::size_t R, std::size_t C, __spv::MatrixLayout U,
+          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
+extern SYCL_EXTERNAL T __spirv_JointMatrixGetSliceElem(
+    __spv::__spirv_JointMatrixINTEL<T, R, C, U, S> *, size_t i);


The width of size_t type is platform dependent I think. Typically it becomes i32 or i64 in LLVM IR. I would prefer a more specific type, like uint32_t.
Is 32 bits enough for indexing in slices, what do you think @dkhaldi?

32 bits is enough. But size_t is widely used in other APIs.
Specifically, what if the user is calculating this iterator using some WI id or other id which is also size_t, will that work?

If we explicitly say that this parameter must be uint32_t there may be a warning about type narrowing during compilation. But now I think maybe I was wrong and we really should stick to size_t. On LLVM IR level we can handle it in the same way as memcpy (https://llvm.org/docs/LangRef.html#llvm-memcpy-intrinsic), i.e. it is overloaded type.

dkhaldi · 2021-12-20T15:50:06Z

sycl/test/matrix/matrix-elemwise-ops.cpp

+             sub_c = joint_matrix_mad(sg, sub_a, sub_b, sub_c);
+             auto wi_slice_c = sub_c.get_wi_data();
+             for (int i = 0; i < wi_slice_c.length(); i++) {
+               wi_slice_c[i] *= 1;


choose a different number other than the neutral element :)

dkhaldi

LGTM

yubingex007-a11y · 2021-12-21T15:51:13Z

ping? @intel/llvm-reviewers-runtime

yubingex007-a11y · 2021-12-22T07:57:49Z

@bader ping? could we merge it?

* upstream/sycl: (1382 commits) [SYCL][XPTI] Report memory allocation info from SYCL runtime (intel#5172) [CI] Switch labels for OCL x64 job (intel#5185) [SYCL] Add basic support for the generic_space address space (intel#5148) [CI] Update CODEOWNERS for SYCL printf support passes (intel#5199) [SYCL][Matrix] Enable wi_slice for joint_matrix (intel#4979) [SYCL][Group algorithms] Move group sort extension to experimental (intel#5169) [SYCL] Fix kernel bundles don't really carry kernel IDs (intel#5121) [SYCL] Initial printf support for non-constant AS format strings (intel#5069) [SYCL][NFC] Fix static code analysis concerns (intel#5189) [SYCL][Doc] Fix typos to fix doc build (intel#5190) [Driver][SYCL] Turn on -fsycl-dead-args-optimization by default (intel#3004) [SYCL][L0][Plugin] Add support for batching copy commands (intel#5155) [CI] Add cache checkout script to docker containers (intel#5184) [SYCL][Doc] Add HIP backend to the filter selector (intel#5176) [Doc] Add documentation for Docker images (intel#4778) [LIBCLC] Add functionality for in-kernel asserts for CUDA backend (intel#5174) Force opt to use new pass manager in exponential-deferred-inlining test after a8c2ba1 [SYCL] Add vec and marray support to known_identity type trait (intel#5163) Correctly resolve merge conflicts Update SPV_INTEL_hw_thread_queries to rev 2 ...

bader · 2021-12-23T08:06:22Z

/summary:run

AlexeySotkin · 2021-12-23T08:11:59Z

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

+                        PI_INVALID_DEVICE);
+#endif // __SYCL_DEVICE_ONLY__
+  }
+  // TODO: add other arithmetic operators


@yubingex007-a11y please do not forget to add overloading for other operators

[Matrix] Enable wi_slice for joint_matrix

16c0bbe

yubingex007-a11y requested a review from a team as a code owner November 17, 2021 16:43

yubingex007-a11y requested review from v-klochkov, dkhaldi and AlexeySotkin November 17, 2021 16:43

yubingex007-a11y commented Nov 17, 2021

View reviewed changes

dkhaldi requested a review from mbelicki November 17, 2021 18:58

dkhaldi reviewed Nov 17, 2021

View reviewed changes

yubingex007-a11y requested a review from LuoYuanke November 18, 2021 02:10

Change implementation by using wi_elem

4803165

dkhaldi reviewed Nov 24, 2021

View reviewed changes

Fix preCI's fail and address dounia's comments

9a867ad

yubingex007-a11y requested a review from MrSidims November 29, 2021 15:23

Remove useless wi_slice_t

d236a56

dkhaldi reviewed Nov 29, 2021

View reviewed changes

Remove useless comments

a894d04

AlexeySotkin reviewed Dec 6, 2021

View reviewed changes

dkhaldi reviewed Dec 14, 2021

View reviewed changes

AlexeySotkin reviewed Dec 15, 2021

View reviewed changes

dkhaldi reviewed Dec 20, 2021

View reviewed changes

Address Dounia&Alexey's comments

93c3e22

yubingex007-a11y added 2 commits December 21, 2021 22:55

Fix clang-format issue

3ffe959

choose a different number for elemwice multiplication

1d2dab3

dkhaldi approved these changes Dec 21, 2021

View reviewed changes

yubingex007-a11y requested review from bader and removed request for AlexeySotkin December 21, 2021 15:50

vladimirlaz approved these changes Dec 21, 2021

View reviewed changes

bader changed the title ~~[Matrix] Enable wi_slice for joint_matrix~~ [SYCL][Matrix] Enable wi_slice for joint_matrix Dec 22, 2021

bader merged commit 97127eb into intel:sycl Dec 22, 2021

AlexeySotkin reviewed Dec 23, 2021

View reviewed changes

dkhaldi mentioned this pull request Jan 7, 2022

[SYCL][Matrix] test the two features: fill a matrix and element wise operations intel/llvm-test-suite#645

Merged

	M.spvm = __spirv_JointMatrixSetSliceElem(M.spvm, idx, rhs);
	M.spvm = __spirv_VectorInsertDynamic(M.spvm, rhs, idx);

		M.spvm = __spirv_JointMatrixSetSliceElem(
		M.spvm, idx, __spirv_JointMatrixGetSliceElem(M.spvm, idx) * rhs);

	return __spirv_JointMatrixGetSliceLength(M.spvm);
	return __spirv_JointMatrixWorkItemLengthINTEL(M.spvm);

	extern SYCL_EXTERNAL T __spirv_JointMatrixGetSliceElem(
	extern SYCL_EXTERNAL T __spirv_VectorExtractDynamic(

		__spirv_JointMatrixSetSliceElem(
		__spv::__spirv_JointMatrixINTEL<T, R, C, U, S> *, size_t i, T val);

[SYCL][Matrix] Enable wi_slice for joint_matrix #4979

[SYCL][Matrix] Enable wi_slice for joint_matrix #4979

Uh oh!

Conversation

yubingex007-a11y commented Nov 17, 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!

Choose a reason for hiding this comment

Uh oh!

yubingex007-a11y commented Dec 14, 2021

Uh oh!

dkhaldi Dec 14, 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!

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!

yubingex007-a11y commented Dec 21, 2021

Uh oh!

yubingex007-a11y commented Dec 22, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

bader commented Dec 23, 2021

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

dkhaldi Dec 14, 2021 •

edited

Loading

yubingex007-a11y commented Dec 22, 2021 •

edited

Loading