- Remove the general query from TODO list as an example is added to t…

…he llvm-test-suite - Add get coord API and remove it from TODO list - Remove the local memory future API looking as it is no more relevant
intel · dm-vodopyanov · Aug 28, 2023 · Jan 9, 2023 · Jan 9, 2023 · Jan 10, 2023
commit 9628bd01b8ff255b8e3adef9e1f2d242dcef4fdb
@@ -151,5 +151,54 @@ The table below provides a list of the combinations that `joint_matrix` implemen
 |  bf16       |  bf16   |   fp32   |  +<=+ 8 |  16   |  16
 |======================
 
+
+## WI data to joint matrix mapping coordinates information for piece-wise operations
+The indexing provided inside the `wi_data` class accesses only the portion of the matrix held by the current WI. It is not possible to know the location of this portion in the joint matrix.  This coordinates mapping  is implementation defined and changes from one backend to the other. For general piece-wise operations like sum of rows of a matrix, the WI data to joint matrix mapping information is needed to reason about the matrix view.
+Within the joint matrix extension, we want to write, as much as possible, one code to run on different backends. If backend X states that a WI owns one exact row of the matrix for instance, writing the following code will work only on that backend for that version of hardware. If a different implementation is used, the same WI may own only half of the row if, for example, the SG size increased.
+The following code locally calculates sum of rows of matrix `joint_matrix<sub_group, int8_t, use::a, 8, 8, layout::row_major> tA;`. In this example, we assume each WI owns 1 columns of `tA` and the sub-group size is 8. `data.length` returns 8 elements per WI.
+
+[frame="none",options="header"]
+|======================
+| a00 | a01 | a02 |a03 | a04 | a05|a06|a07|a08|a09|a010|a011|a012|a013|a014|a015
+| a10 | a11 | a12 |a13 | a14 | a15|a16|a17|a18|a19|a110|a111|a112|a113|a114|a115
+| a20 | a21 | a22 |a23 | a24 | a25|a26|a27|a28|a29|a210|a211|a212|a213|a214|a215
+| a30 | a31 | a32 |a33 | a34 | a35|a36|a37|a38|a39|a310|a311|a312|a313|a314|a315
+| a40 | a41 | a42 |a43 | a44 | a45|a46|a47|a48|a49|a410|a411|a412|a413|a414|a415
+| a50 | a51 | a52 |a53 | a54 | a55|a56|a57|a58|a59|a510|a511|a512|a513|a514|a515
+| a60 | a61 | a62 |a63 | a64 | a65|a66|a67|a68|a69|a610|a611|a612|a613|a614|a615
+| a70 | a71 | a72 |a73 | a74 | a75|a76|a77|a78|a79|a710|a711|a712|a713|a714|a715
+|======================
+
+
+```c++
+auto data = get_wi_data(sg, tA);
+// each WI calculates local sum of rows
+for (int row = 0; row < 8; row++) { 
+  for (int i = 0; i < data.length()/8; ++i) {//WI owns 1 element per row
+    sum_of_local_rows[row] += data[i+row];
+  }
+}  
+```
+The code above assumes knowledge of the distribution of the joint matrix across the different work-items. This is different when a different distribution happens. In order to be agnostic to this distribution, instead of hard-coding this mapping, we use general backend and target-agnostic functionality. To this end, a new method is added to 'wi_element' to query this mapping.
+
+```c++
+namespace sycl::ext::intel::experimental::matrix {
+ std::tuple<uint32_t, uint32_t> get_coord();
+}
+```
+
+`get_coord` returns [row,col] coordinates of the current object `wi_element` of the joint matrix.  The code above results into the following:
+
+
+```c++
+auto data = get_wi_data(sg, tA);
+// each WI calculates local sum of rows
+for (int i = 0; i < data.length(); ++i) {
+  auto [row, col] = data[i].get_coord();
+  sum_of_local_rows[row] += data[i];
+}  
+```
+
+
 ## Open Questions
 - Should the same class, `joint_matrix`, handle both cases where sizes are constant (GPU case) and when sizes are variable (CPU case)? Note that a Intel AMX 2d tile register permits sizes up to 1024 (16rowsx64cols) bytes that can be variable. The ability to define only one interface for both would make it possible to give the user a way to make use of the flexibility introduced by the CPU but at the same time save resources on the GPU. In a previous version of the design, we used `sycl::dynamic_extent`  to differentiate between static and dynamic sizes. But since this was not implemented at all, we decided to remove it. We can revisit this design choice if this comes up as part of a customer request or if SPIRV matrix extension extends its support to dynamic sizes.
diff --git a/.../extensions/experimental/sycl_ext_oneapi_matrix/sycl_ext_oneapi_matrix.asciidoc b/.../extensions/experimental/sycl_ext_oneapi_matrix/sycl_ext_oneapi_matrix.asciidoc
@@ -546,43 +546,6 @@ joint_matrix<sub_group, int, use::accumulator, msize, nsize> sub_c;
 //Remainder handling
 ```
 
-## Future-looking API
-
-### Memory scope
-The current experimental API uses `joint_` semantics to define the memory scope of the matrix. The long term solution is to use the proposed link:../supported/sycl_ext_oneapi_local_memory.asciidoc[`group_local_memory` extension] to allocate the matrix in local memory associated with a SYCL group as shown in the example below.
-
-
-```c++
-multi_ptr<matrix<T>, address_space::local_space> tA_ptr = group_local_memory<matrix<sub_group, int8_t, tM, tN, use::a>>(sg);
-```
-We did not utilize this extension for this matrix API version because sub-group local memory is not yet well defined in {dpcpp}. Moreover, the representation of this notion in LLVM IR and SPIR-V is not clear yet. 
-
-### WI data to joint matrix mapping coordinates information for piece-wise operations
-The indexing provided inside the `wi_data` class accesses only the portion of the matrix held by the current WI. It is not possible to know the location of this portion in the original matrix.  This coordinates mapping  is implementation defined and changes from one backend to the other. For general piece-wise operations like sum of rows of a matrix, the WI data to joint matrix mapping information is needed to reason about the matrix view.
-Within the joint matrix extension, we want to write, as much as possible, one code to run on different backends. If backend X states that a WI owns one exact row of the matrix for instance, writing the following code will work only on that backend for that version of hardware. If a different hardware and implementation is used, the same WI may own only half of the row if, for example, the SG size increased. 
-
-```c++
-auto data = get_wi_data(sg, C);
-for (int i = 0; i < data.length(); ++i) {
-  sum_of_local_rows[row] += data[i];
-}
-```
-
-We want to keep backward compatibility in the joint matrix code when implementations or hardware change. To that end, instead of hard-coding this mapping, we use general backend and target-agnostic functionality, especially in the JIT compilation mode of SYCL. For this reason we would like to be able to query this mapping so that code does not have to change from one version to the other.
-
-So for the mapping problem, since this mapping is implementation-defined, one of the proposals is to add runtime functions like:
-```c++
-auto data = get_wi_data(sg, C);
-for (int i = 0; i < data.length; ++i) {
-  auto row, col = data[i].get_coord();
-  sum_of_local_rows[row] += data[i];
-}
-```
-
-## TODO List
-- Add WI data to joint matrix mapping coordinates information for piece-wise operations. This will be added as part of the query or new methods to the 'get_wi_data' class. 
-- Add a more realistic and complete example that shows the value of the general query. 
-
 
 ## Revision History