Skip to content

Commit c64b157

Browse files
dkhaldidkhaldi
authored
[SYCL] Move the Intel specific features to a separate document (#7307)
 Mainly: - Remove dynamic extent. I kept it in the open questions though of the Intel specific document. - Remove the extra functionality Intel implementation provides WRT layout to the Intel specific document Besides that, I also change the name of DPAS to Intel XMX in this PR
1 parent b2d6fdf commit c64b157

File tree

2 files changed

+249
-160
lines changed

2 files changed

+249
-160
lines changed

‎sycl/doc/extensions/experimental/sycl_ext_oneapi_matrix/sycl_ext_intel_matrix.asciidoc

Lines changed: 155 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,155 @@
1+
# Additional Intel-only specifics about matrix extension for DPC++
2+
3+
:source-highlighter: coderay
4+
:coderay-linenums-mode: table
5+
:dpcpp: pass:[DPC++]
6+
7+
// This section needs to be after the document title.
8+
:doctype: book
9+
:toc2:
10+
:toc: left
11+
:encoding: utf-8
12+
:lang: en
13+
14+
:blank: pass:[ +]
15+
16+
// Set the default source code type in this document to C++,
17+
// for syntax highlighting purposes. This is needed because
18+
// docbook uses c++ and html5 uses cpp.
19+
:language: {basebackend@docbook:c++:cpp}
20+
21+
22+
== Notice
23+
24+
Copyright (c) 2021-2022 Intel Corporation. All rights reserved.
25+
26+
NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are
27+
trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc.
28+
used by permission by Khronos.
29+
30+
This extension is written against the SYCL 2020 revision 5 specification. All
31+
references below to the "core SYCL specification" or to section numbers in the
32+
SYCL specification refer to that revision.
33+
34+
**_NOTE:_** This document describes the extra features and details for the implementation of `joint_matrix` extension on Intel AMX and Intel XMX.
35+
This is an initial experimental version to try out functionality
36+
and performance, and **future versions of this API may change in ways that are incompatible with this experimental version**.
37+
38+
## Introduction
39+
The Intel backend implementations on both Intel AMX and Intel XMX support `joint_matrix`, `joint_matrix_load`, `joint_matrix_store`, `joint_matrix_mad`, `joint_matrix_fill`, `get_wi_data`, and the query interface, as they are defined in the sycl_ext_oneapi_matrix extension. There are additional specifics about the supported layouts that enable extra performance and functionality listed in this document.
40+
This extension presents some supplementary Intel AMX and Intel XMX features not contained within the sycl_ext_oneapi_matrix extension. The additional features are built on top of the sycl_ext_oneapi_matrix extension but are only supported by the Intel AMX and Intel XMX backends.
41+
42+
## Feature test macro
43+
44+
This extension provides a feature-test macro as described in the core SYCL
45+
specification section 6.3.3 "Feature test macros". Therefore, an
46+
implementation supporting this extension must predefine the macro
47+
`SYCL_EXT_INTEL_MATRIX` to one of the values defined in the table below.
48+
Applications can test for the existence of this macro to determine if the
49+
implementation supports this feature, or applications can test the macro's
50+
value to determine which of the extension's APIs the implementation supports.
51+
52+
[frame="none",options="header"]
53+
|======================
54+
|Value |Description
55+
|1 |Introduce `packed` layout and extend `joint_matrix_store` to Matrix A and B.
56+
|======================
57+
58+
59+
## Extra Functionality
60+
61+
### Layout
62+
Besides row major and column major layouts, `layout` introduces the custom layout packed layout that refers to the VNNI format descibed in the following section.
63+
64+
```c++
65+
namespace sycl::ext::intel::experimental::matrix {
66+
enum class layout {
67+
packed
68+
};
69+
}
70+
```
71+
72+
73+
### Layout argument in `joint_matrix_load`
74+
`layout` in `joint_matrix_load` can take `packed` as argument to specify that the data has already been transformed into VNNI format (`packed`). in this case, `stride` argument of `joint_matrix_load` describes the number of elements between consecutive rows for packed layouts.
75+
76+
In order to get maximum performance on Intel AMX and Intel XMX, prepacking data in the memory is necessary. If users did not specify the packed layouts, transforms done by the implementation will be slow due to extra scatter/gather operations. Hence, we expose the `packed` layout to the user to specify that A or B have already been VNNIed. The packed or VNNI layout is introduced in the `VNNI layout` section below.
77+
78+
IMPORTANT: In the current Intel AMX and Intel XMX implementations, the layout in the load of matrix B (provided by the `layout memL` parameter below) must be `packed` or `row_major`. Automatic VNNI transform is supported on AMX. The layout in the load of matrices A and C must be `row_major`, and the layout in the store of matrix C (provided by the `layout memL` parameter below) must also be `row_major`.
79+
80+
### Store Operation
81+
Besides store of matrix `accumulator`, the Intel implementation allows store on matrix `a` and `b` as well.
82+
83+
#### Store
84+
```c++
85+
namespace sycl::ext::intel::experimental::matrix {
86+
template <typename Group, typename T, size_t NumRows, size_t NumCols,
87+
use Use, layout Layout, access::address_space Space>
88+
void joint_matrix_store(Group sg,
89+
joint_matrix<Group, T, Use, NumRows, NumCols, Layout> &res,
90+
multi_ptr<T, Space, IsDecorated> src, size_t stride);
91+
}
92+
```
93+
94+
95+
## VNNI/Packed Layout
96+
Intel AMX and Intel XMX compute assumes that the B tile register (src1) is in the VNNI format as they need 32bit of K-data in A and B to be contiguous in memory.
97+
The VNNI blocking factor is 2 in the case of 16-bit types, and it is 4 in the case of 8-bit types. While the current implementation assumes that the matrix has been already packed by the user for performance reasons, the layout information is needed to inform the implementation about this transformation. The following example illustrates how a matrix in `row_major` layout is transformed into the `packed` layout for a 16-bit type.
98+
99+
#### Example 1: 16-bit elements
100+
// Example of a 4 row x 4 column matrix using a 16-bit data element, in row-major layout.
101+
// Element a1 is contiguous in memory with element b1, etc.
102+
// ---------------------------------
103+
// a1, b1, c1, d1
104+
// a2, b2, c2, d2
105+
// a3, b3, c3, d3
106+
// a4, b4, c4, d4
107+
// ---------------------------------
108+
// The same matrix reformatted in packed layout.
109+
// Here, packing of 2 elements is needed to form 32 bits.
110+
// Element a1 is contiguous in memory with element a2, etc.
111+
// ---------------------------------
112+
// a1, a2, b1, b2, c1, c2, d1, d2
113+
// a3, a4, b3, b4, c3, c4, d3, d4
114+
115+
#### Example 2: 8-bit elements
116+
117+
// Example of a 4 row x 4 column matrix using a 8-bit data element, in row-major layout.
118+
// Element a1 is contiguous in memory with element b1, etc.
119+
// ---------------------------------
120+
// a1, b1, c1, d1
121+
// a2, b2, c2, d2
122+
// a3, b3, c3, d3
123+
// a4, b4, c4, d4
124+
// ---------------------------------
125+
// The same matrix reformatted in packed layout.
126+
// Here, packing of 4 elements is needed to form 32 bits.
127+
// Elements a1, a2, a3, a4 are contiguous in memory, etc.
128+
// ---------------------------------
129+
// a1, a2, a3, a4, b1, b2, b3, b4, c1, c2, c3, c4, d1, d2, d3, d4
130+
131+
## Supported Combinations Per Hardware
132+
133+
The table below provides a list of the combinations that `joint_matrix` implementations support on each of Intel AMX and Intel XMX hardware. Note that these can be returned in a parametrized way using the `tpu_params` query class.
134+
135+
### Intel AMX Supported Combinations
136+
137+
[frame="none",options="header"]
138+
|======================
139+
| A type | B type | Accumulator type | M | N | K
140+
| (u)int8_t | (u)int8_t | int32_t | +<=+ 16 | +<=+ 16 | +<=+ 64
141+
| bf16 | bf16 | fp32 | +<=+ 16 | +<=+ 16 | +<=+ 32
142+
|======================
143+
144+
### Intel XMX Supported Combinations
145+
146+
[frame="none",options="header"]
147+
|======================
148+
| A type | B type | Accumulator type | M | N | K
149+
| (u)int8_t | (u)int8_t | int32_t | +<=+ 8 | 16 | 32
150+
| fp16 | fp16 | fp32 | +<=+ 8 | 16 | 16
151+
| bf16 | bf16 | fp32 | +<=+ 8 | 16 | 16
152+
|======================
153+
154+
## Open Questions
155+
- 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.

0 commit comments

Comments
 (0)