Skip to content

[SYCL][Docs][Joint Matrix] Add more devices support and more types and shape combinations #15547

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Oct 21, 2024
Merged

[SYCL][Docs][Joint Matrix] Add more devices support and more types and shape combinations #15547

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
528581b
[SYCL][Docs][Joint Matrix] Add more types and shape combinations for PVC
dkhaldi Sep 28, 2024
a4b9270
Add Battlemage, Lunar Lake and Arrow Lake H to the combinations appendix
dkhaldi Oct 15, 2024
197d680
fix a typo
dkhaldi Oct 15, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion sycl/doc/extensions/experimental/sycl_ext_matrix/sycl_ext_intel_matrix.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -113,7 +113,7 @@ an application can use to indicate that the matrix data is loaded or
stored in VNNI "packed" format.

```c++
namespace sycl::ext::oneapi::experimental::matrix::layout {
namespace sycl::ext::oneapi::experimental::matrix {

enum class layout {
ext_intel_packed
Expand Down
113 changes: 82 additions & 31 deletions sycl/doc/extensions/experimental/sycl_ext_matrix/sycl_ext_oneapi_matrix.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1017,52 +1017,103 @@ is shown in a single column in the table below.

==== Intel XMX Supported Combinations
This is currently available in devices with the architecture
`architecture::intel_gpu_pvc`, `architecture::intel_gpu_dg2_g10`,
`architecture::intel_gpu_dg2_g11`, and
`architecture::intel_gpu_dg2_g12`.
In these architectures'
implementation, the type of the C matrix must be the same as the type
of the D matrix. Therefore, that common type is shown in a single
column in the table below.
`architecture::intel_gpu_pvc`, `architecture::intel_gpu_bmg_g21`,
`architecture::intel_gpu_lnl_m`, `architecture::intel_gpu_dg2_g10`,
`architecture::intel_gpu_dg2_g11`, `architecture::intel_gpu_dg2_g12`,
and `architecture::intel_gpu_arl_h`.

[frame="none",options="header"]
|======================
| A type | B type | C and D type | M | N | K | device
| A type | B type | C type | D type | M | N | K | device
.2+| `matrix_type::uint8` .2+| `matrix_type::uint8` .2+|
`matrix_type::sint32` .2+| +<=+ 8 | 16 .2+| 32
|`architecture::intel_gpu_pvc`
`matrix_type::sint32` .2+| `matrix_type::sint32` .2+| +<=+ 8 | 16 .2+| 32
|`architecture::intel_gpu_pvc`, `architecture::intel_gpu_bmg_g21`,
`architecture::intel_gpu_lnl_m`
|8|`architecture::intel_gpu_dg2_g10,
architecture::intel_gpu_dg2_g11, architecture::intel_gpu_dg2_g12`
architecture::intel_gpu_dg2_g11, architecture::intel_gpu_dg2_g12`,
`architecture::intel_gpu_arl_h`
.2+| `matrix_type::uint8` .2+| `matrix_type::sint8` .2+|
`matrix_type::sint32` .2+| +<=+ 8 | 16 .2+| 32 |
`architecture::intel_gpu_pvc`
`matrix_type::sint32` .2+|`matrix_type::sint32` .2+| +<=+ 8 | 16 .2+| 32 |
`architecture::intel_gpu_pvc`, `architecture::intel_gpu_bmg_g21`,
`architecture::intel_gpu_lnl_m`
|8|`architecture::intel_gpu_dg2_g10,
architecture::intel_gpu_dg2_g11, architecture::intel_gpu_dg2_g12`
architecture::intel_gpu_dg2_g11, architecture::intel_gpu_dg2_g12`,
`architecture::intel_gpu_arl_h`
.2+| `matrix_type::sint8` .2+| `matrix_type::uint8` .2+|
`matrix_type::sint32` .2+| +<=+ 8 | 16 .2+| 32 |
`architecture::intel_gpu_pvc`
`matrix_type::sint32` .2+|`matrix_type::sint32` .2+| +<=+ 8 | 16 .2+| 32 |
`architecture::intel_gpu_pvc`, `architecture::intel_gpu_bmg_g21`,
`architecture::intel_gpu_lnl_m`
|8|`architecture::intel_gpu_dg2_g10,
architecture::intel_gpu_dg2_g11, architecture::intel_gpu_dg2_g12`
architecture::intel_gpu_dg2_g11, architecture::intel_gpu_dg2_g12`,
`architecture::intel_gpu_arl_h`
.2+| `matrix_type::sint8` .2+| `matrix_type::sint8` .2+|
`matrix_type::sint32` .2+| +<=+ 8 | 16 .2+| 32 |
`architecture::intel_gpu_pvc`
`matrix_type::sint32` .2+| `matrix_type::sint32` .2+| +<=+ 8 | 16 .2+| 32 |
`architecture::intel_gpu_pvc`, `architecture::intel_gpu_bmg_g21`,
`architecture::intel_gpu_lnl_m`
|8|`architecture::intel_gpu_dg2_g10,
architecture::intel_gpu_dg2_g11, architecture::intel_gpu_dg2_g12`
.2+|`matrix_type::fp16` .2+| `matrix_type::fp16` .2+|
`matrix_type::fp32` .2+| +<=+ 8 | 16 .2+| 16 |
`architecture::intel_gpu_pvc`
|8| `architecture::intel_gpu_dg2_g10,
architecture::intel_gpu_dg2_g11, architecture::intel_gpu_dg2_g12`
.6+| `matrix_type::bf16` .6+| `matrix_type::bf16` .6+|
`matrix_type::fp32` | 16 | 16 | 16 .4+|`architecture::intel_gpu_pvc`
| 1 | 64 | 16 | 32 | 64 | 16
architecture::intel_gpu_dg2_g11, architecture::intel_gpu_dg2_g12`,
`architecture::intel_gpu_arl_h`
.8+|`matrix_type::fp16` .8+| `matrix_type::fp16` .8+|
`matrix_type::fp32` .8+|`matrix_type::fp32` .1+| 16 .1+| 16 | 16
.6+|`architecture::intel_gpu_pvc`, `architecture::intel_gpu_bmg_g21`,
`architecture::intel_gpu_lnl_m`
.2+| 1 .2+| 64 | 16 |32
.2+| 32 .2+| 64 | 16 |32
.2+| +<=+ 8 | 16 .2+| 16
|8 .2+| `architecture::intel_gpu_dg2_g10,
architecture::intel_gpu_dg2_g11, architecture::intel_gpu_dg2_g12`,
`architecture::intel_gpu_arl_h`
.1+| 32 .1+| 32 .1+| 16
.6+|`matrix_type::fp16` .6+| `matrix_type::fp16` .6+|
`matrix_type::fp16` .6+|`matrix_type::fp32` .1+| +<=+ 8 | 16 .1+| 16
.6+| `architecture::intel_gpu_pvc`, `architecture::intel_gpu_bmg_g21`,
`architecture::intel_gpu_lnl_m`
| 16 | 16 | 16 .2+| 1 .2+| 64 | 16 | 32
.2+| 32 .2+| 64 | 16 | 32
.6+|`matrix_type::fp16` .6+| `matrix_type::fp16` .6+|
`matrix_type::fp32` .6+|`matrix_type::fp16` .1+| +<=+ 8 | 16 .1+| 16
.6+|`architecture::intel_gpu_pvc`, `architecture::intel_gpu_bmg_g21`,
`architecture::intel_gpu_lnl_m`
| 16 | 16 | 16 .2+| 1 .2+| 64 | 16 | 32
.2+| 32 .2+| 64 |16 | 32
.6+|`matrix_type::fp16` .6+| `matrix_type::fp16` .6+|
`matrix_type::fp16` .6+|`matrix_type::fp16` .1+| +<=+ 8 | 16 .1+| 16
.6+|`architecture::intel_gpu_pvc`, `architecture::intel_gpu_bmg_g21`,
`architecture::intel_gpu_lnl_m`
| 16 | 16 | 16 .2+| 1 .2+| 64 | 16 |32 .2+| 32 .2+| 64 | 16 | 32
.8+| `matrix_type::bf16` .8+| `matrix_type::bf16` .8+|
`matrix_type::fp32` .8+| `matrix_type::fp32` | 16 | 16 | 16
.6+|`architecture::intel_gpu_pvc`, `architecture::intel_gpu_bmg_g21`,
`architecture::intel_gpu_lnl_m`
.2+| 1 .2+| 64 | 16 | 32
.2+| 32 .2+| 64 | 16 |32
.2+| +<=+ 8 | 16 .2+| 16
|8 .2+| `architecture::intel_gpu_dg2_g10,
architecture::intel_gpu_dg2_g11, architecture::intel_gpu_dg2_g12`
architecture::intel_gpu_dg2_g11, architecture::intel_gpu_dg2_g12`,
`architecture::intel_gpu_arl_h`
.1+| 32 .1+| 32 .1+| 16
.6+|`matrix_type::bf16` .6+| `matrix_type::bf16` .6+|
`matrix_type::bf16` .6+|`matrix_type::fp32` .1+| +<=+ 8 | 16 .1+| 16 .6+|
`architecture::intel_gpu_pvc`, `architecture::intel_gpu_bmg_g21`,
`architecture::intel_gpu_lnl_m`
| 16 | 16 | 16 .2+| 1 .2+| 64 | 16 | 32
.2+| 32 .2+| 64 |16 | 32
.6+|`matrix_type::bf16` .6+| `matrix_type::bf16` .6+|
`matrix_type::fp32` .6+|`matrix_type::bf16` .1+| +<=+ 8 | 16 .1+| 16 .6+|
`architecture::intel_gpu_pvc`, `architecture::intel_gpu_bmg_g21`,
`architecture::intel_gpu_lnl_m`
| 16 | 16 | 16 .2+| 1 .2+| 64 | 16 | 32
.2+| 32 .2+| 64 |16 | 32
.6+|`matrix_type::bf16` .6+| `matrix_type::bf16` .6+|
`matrix_type::bf16` .6+|`matrix_type::bf16` .1+| +<=+ 8 | 16 .1+| 16 .6+|
`architecture::intel_gpu_pvc`, `architecture::intel_gpu_bmg_g21`,
`architecture::intel_gpu_lnl_m`
| 16 | 16 | 16 .2+| 1 .2+| 64 | 16 | 32
.2+| 32 .2+| 64 |16 | 32
| `matrix_type::tf32` | `matrix_type::tf32` |
`matrix_type::fp32` | +<=+ 8 | 16 | 8 |
`architecture::intel_gpu_pvc`
`matrix_type::fp32` .2+| `matrix_type::fp32` | +<=+ 8 | 16 | 8 |
`architecture::intel_gpu_pvc`, `architecture::intel_gpu_bmg_g21`,
`architecture::intel_gpu_lnl_m`
|======================

==== Nvidia Tensor Cores Supported Combinations
Expand Down