Misc. bug: XMX cores not detected on Intel Arc B580

[codecnotsupported]

Name and Version

version: 1 (eb7cf15)
built with cc (GCC) 14.2.1 20240910 for x86_64-pc-linux-gnu

Operating systems

Linux

Which llama.cpp modules do you know to be affected?

Other (Please specify in the next section)

Command line

`$ ./build/bin/llama-ls-sycl-device`

Found 1 SYCL devices:
|  |                   |                                       |       |Max    |        |Max  |Global |                     |         XMX  |
|  |                   |                                       |       |compute|Max work|sub  |mem    |                     |          or  |
|ID|        Device Type|                                   Name|Version|units  |group   |group|size   |       Driver version| Tensor Cores |
|--|-------------------|---------------------------------------|-------|-------|--------|-----|-------|---------------------|--------------|
| 0| [level_zero:gpu:0]|                Intel Graphics [0xe20b]|   20.1|    160|    1024|   32| 12168M|            1.6.32224|            no|

XMX or Tensor Cores: no
But the B580 does have XmX cores.

Problem description & steps to reproduce

N/A

First Bad Commit

Unknown

Relevant log output

N/A

[codecnotsupported]

It should be already supported in SYCL as of this commit.
intel/llvm@4bcea94

[qnixsynapse]

Hey! Thanks for the issue.

But I don't think this is a problem with ggml sycl backend, because we are using this to determine whether a device supports intel matrix or not. Please see:

https://github.com/ggerganov/llama.cpp/blob/eb7cf15a808d4d7a71eef89cc6a9b96fe82989dc/ggml/src/ggml-sycl/common.cpp#L54-L56

First thing is to verify whether that commit is in the release builds of Intel's LLVM/oneapi DPC++ compiler.

For example, here is what it shows for Arc A750:

Found 1 SYCL devices:
|  |                   |                                       |       |Max    |        |Max  |Global |                     |         XMX  |
|  |                   |                                       |       |compute|Max work|sub  |mem    |                     |          or  |
|ID|        Device Type|                                   Name|Version|units  |group   |group|size   |       Driver version| Tensor Cores |
|--|-------------------|---------------------------------------|-------|-------|--------|-----|-------|---------------------|--------------|
| 0| [level_zero:gpu:0]|                Intel Arc A750 Graphics|  12.55|    448|    1024|   32|  8096M|            1.6.32224|           yes|

[codecnotsupported]

Yep, not a issue with llama.cpp, the Intel® oneAPI Base Toolkit has a dated version of oneAPI DPC++.
Which doesn't advertise sycl::aspect::ext_intel_matrix.

Fun fact, ~8 days after the last stable release of oneAPI DPC++, there was a PR to advertise sycl::aspect::ext_intel_matrix for the B580, guess I got a bit unlucky there with the timing.

Anyhow closing for now, should be fixed with the latest oneAPI DPC++ release. Or when the Intel® oneAPI Base Toolkit update releases.

[NeoZhangJianyu]

@qnixsynapse
XMX is not needed to be shown in this table.
It's not important now, because this feature is not used in most cases.

"Tensor Core" is the CUDA concept, I don't think the API: sycl::aspect::ext_intel_matrix support to detect "Tensor Core" in CUDA device.

I suggest hiding this column (XMX or Tensor Core) in the table, to avoid misunderstanding.

It's OK to keep the function: gpu_has_xmx().

[qnixsynapse]

@NeoZhangJianyu Thanks. Will cover this in #11515

[NeoZhangJianyu]

OK, thank you!

[codecnotsupported]

@qnixsynapse XMX is not needed to be shown in this table. It's not important now, because this feature is not used in most cases.

"Tensor Core" is the CUDA concept, I don't think the API: sycl::aspect::ext_intel_matrix support to detect "Tensor Core" in CUDA device.

I suggest hiding this column (XMX or Tensor Core) in the table, to avoid misunderstanding.

It's OK to keep the function: gpu_has_xmx().

You want to hide the XMX or Tensor Core column because sycl::aspect::ext_intel_matrix doesn't detect Cuda Tensor cores.
But XMX is still correctly detected/supported right? Wouldn't it be better to simply rename it to XMX or Tensor Core to XMX?
If XMX provides a speedup it would be useful to know if it is supported.

[Alcpz]

Yep, not a issue with llama.cpp, the Intel® oneAPI Base Toolkit has a dated version of oneAPI DPC++. Which doesn't advertise sycl::aspect::ext_intel_matrix.

Fun fact, ~8 days after the last stable release of oneAPI DPC++, there was a PR to advertise sycl::aspect::ext_intel_matrix for the B580, guess I got a bit unlucky there with the timing.

Anyhow closing for now, should be fixed with the latest oneAPI DPC++ release. Or when the Intel® oneAPI Base Toolkit update releases.

It is fixed. See the output using the open source version of DPCPP

|  |                   |                                       |       |Max    |        |Max  |Global |                     |         XMX  |                                                                                           
|  |                   |                                       |       |compute|Max work|sub  |mem    |                     |          or  |                                                                                           
|ID|        Device Type|                                   Name|Version|units  |group   |group|size   |       Driver version| Tensor Cores |              
|--|-------------------|---------------------------------------|-------|-------|--------|-----|-------|---------------------|--------------|                                                                                           
| 0| [level_zero:gpu:0]|                Intel Graphics [0xe20b]|   20.1|    160|    1024|   32| 12168M|         1.6.31740+11|           yes|                                                                                           
| 1| [level_zero:gpu:1]|                 Intel UHD Graphics 770|   12.2|     32|     512|   32| 61332M|         1.6.31740+11|            no|

We will need to wait for the next release.