[SYCL][Intel GPU] Long Term Features & Issues Tracking

[airMeng]

Feel free to drop a note, let's know if you have any feature request or bugs (even unconfirmed)

• Multi-card Support. Issue #5282, PR #5806
• Multi-batch Support #5272 Low performance with Sycl Backend #5480
• CI test error for more than one GPU is detected and used.
  Current code returns all SYCL devices, including CPU, GPU (level-zero, opencl), FPGA. SYCL only support GPU. So when CI test on other devices, it will be fault.
• Support no-mmap parameter in other application.
  There is known issue of SYCL: memcpy() from host (mmap) to device will hang in same cases. It's not resolved now. A work around solution is no use mmap. I have handled it in llama-bench (add --mmap parameter). We need add to more applications in examples.
• Clean code for warning and unused macro and variable.
  Suggest to handle it after multiple-card is finished. Lots of such unused code will be useful for multiple-card feature.
• Support SYCL build for Nvidia and AMD targets #5357
• Improve first token performance.

Also let's know if you have taken any tasks here.

cc @NeoZhangJianyu @luoyu-intel @abhilash1910

[NeoZhangJianyu]

I'd like handle: Multi-card Support, CI test error for more than one GPU is detected and used.

[abhilash1910]

For code cleaning and sanitization of compiler runtime I will be adding patch post the previous changes
Also sycl for other vendor build targets.

[DatCaptainHorse]

Would like to see support for SOTA 2-bit quantized models (GGML_TYPE_IQ2_XXS, GGML_TYPE_IQ2_XS, GGML_TYPE_IQ3_XXS).

Been trying to do this myself for past hour or two, using dpct isn't as trouble-free as it makes you think.

[characharm]

before airMeng:sycl_fix_max_alloc_size

I don't know if these numbers can be considered when the model produces useless output.

Using device 0 (Intel(R) Arc(TM) A770 Graphics) as main device

model	size	params	backend	ngl	test	t/s
llama 13B Q5_K - Medium	8.60 GiB	13.02 B	GPU BLAS	99	pp 512	430.52 ± 38.83
llama 13B Q5_K - Medium	8.60 GiB	13.02 B	GPU BLAS	99	tg 128	16.14 ± 0.13

model	size	params	backend	ngl	test	t/s
llama 7B Q6_K	5.53 GiB	7.24 B	GPU BLAS	99	pp 512	734.90 ± 149.16
llama 7B Q6_K	5.53 GiB	7.24 B	GPU BLAS	99	tg 128	22.00 ± 0.17

after

model	size	params	backend	ngl	test	t/s
llama 13B Q5_K - Medium	8.60 GiB	13.02 B	SYCL	99	pp 512	392.43 ± 33.81
llama 13B Q5_K - Medium	8.60 GiB	13.02 B	SYCL	99	tg 128	9.75 ± 0.06

model	size	params	backend	ngl	test	t/s
llama 7B Q6_K	5.53 GiB	7.24 B	SYCL	99	pp 512	703.17 ± 122.82
llama 7B Q6_K	5.53 GiB	7.24 B	SYCL	99	tg 128	19.07 ± 0.15

vulkan latest

model	size	params	backend	ngl	test	t/s
llama 13B Q5_K - Medium	8.60 GiB	13.02 B	Vulkan	99	pp 512	33.99 ± 0.53
llama 13B Q5_K - Medium	8.60 GiB	13.02 B	Vulkan	99	tg 128	8.04 ± 0.01

model	size	params	backend	ngl	test	t/s
llama 7B Q6_K	5.53 GiB	7.24 B	Vulkan	99	pp 512	74.76 ± 3.49
llama 7B Q6_K	5.53 GiB	7.24 B	Vulkan	99	tg 128	21.26 ± 0.10

[qnixsynapse]

I wonder why build fails with -DLLAMA_SYCL_F16=ON for my Intel Arc 750..

I think this GPU support f16.

But anyways getting 16 tokens/sec for q4_k_m for 7B with all layers on GPU.

With batch bench, mitral -7B Q4_K_M:

PP	TG	B	N_KV	T_PP s	S_PP t/s	T_TG s	S_TG t/s	T s	S t/s
128	128	1	256	0.497	257.31	8.129	15.75	8.626	29.68
128	128	2	512	0.539	475.29	43.357	5.90	43.896	11.66
128	128	4	1024	0.799	641.06	30.104	17.01	30.903	33.14
128	128	8	2048	1.546	662.36	30.739	33.31	32.285	63.44
128	128	16	4096	3.261	627.93	32.109	63.78	35.371	115.80
128	256	1	384	0.338	378.75	15.331	16.70	15.669	24.51
128	256	2	768	0.452	566.45	86.797	5.90	87.249	8.80
128	256	4	1536	0.759	674.89	60.522	16.92	61.281	25.06
128	256	8	3072	1.548	661.63	61.703	33.19	63.251	48.57

[NeoZhangJianyu]

@akarshanbiswas
It should be fixed by PR: #5411.
The FP16 build is added to CI. There won't be build issue of FP16 in the feature.

By the way, if you have more issues, please create/update in the "issue" field for it, instead of "discussion" field.

[qnixsynapse]

Oh okay. The title of this discussion has "issues tracking" so I mentioned that issue here. :)

[airMeng]

@akarshanbiswas Could you report the full performance here like the above comments since we don't have A750 on our side, so a performance baseline on A750 would be quite helpful? Thank you for your effort to help improving SYCL backend.

You can find the standard performance measurements here

[qnixsynapse]

@airMeng Thank you. I was able to test a little and I have updated my comment here.

[kunger97]

I would like to know if there are plans to support quantization types that are not currently supported like iq3, iq4?

[characharm]

With the recent changes, the model nous-hermes-2-34b-2.69 has seen a significant speedup. From the unusable 3-4 tokens per second, it now reaches 7-8

detect 1 SYCL GPUs: [0,] with Max compute units:512

model	size	params	backend	ngl	test	t/s
llama 30B Q2_K - Medium	10.76 GiB	34.39 B	SYCL	99	pp 512	159.76 ± 8.53
llama 30B Q2_K - Medium	10.76 GiB	34.39 B	SYCL	99	tg 128	8.52 ± 0.02

build: 21b0867 (2345)

for comparison

Vulkan0: Intel(R) Arc(TM) A770 Graphics | uma: 0 | fp16: 1 | warp size: 32

model	size	params	backend	ngl	test	t/s
llama 30B Q2_K - Medium	10.76 GiB	34.39 B	Vulkan	99	pp 512	33.27 ± 0.29
llama 30B Q2_K - Medium	10.76 GiB	34.39 B	Vulkan	99	tg 128	5.05 ± 0.05

build: 82cb31e (2348)

I hope other quantization methods will also see an improvement, for now, they more or less perform similarly to Vulkan.

[qnixsynapse]

I have a question. During prompt processing or generation, the llama.cpp's SYCL backend seems to use only one of the (I am assuming XMX) engines of my GPU. Although they are tagged 'unknown' in intel_gpu_top. Does anyone knows why? And is it possible to parallelize across all of them?

On Linux, I can't even monitor neither the VRAM usage, nor the temps which is surprising because it is a while since this GPU was launched. My all hopes for the new Xe driver.

[airMeng]

I have a question. During prompt processing or generation, the llama.cpp's SYCL backend seems to use only one of the (I am assuming XMX) engines of my GPU. Although they are tagged 'unknown' in intel_gpu_top. Does anyone knows why? And is it possible to parallelize across all of them?

On Linux, I can't even monitor neither the VRAM usage, nor the temps which is surprising because it is a while since this GPU was launched. My all hopes for the new Xe driver.

@akarshanbiswas can you try https://github.com/intel/xpumanager to monitor the usage?

[qnixsynapse]

It is not yet available for Arch Linux. I think I will have to compile it from source. Also, it is for datacenter GPUs. Will it support Intel Arc?

p.s I didn't find any sysfs interface for monitoring temps or vram usage. Will this work with currently stable mainline kernel?

[airMeng]

will it support Intel Arc

yes it is.

I didn't find any sysfs interface for monitoring temps or vram usage

you can get temps and vram usage here https://github.com/intel/xpumanager/blob/fdcb817c0dfadf9423da4dfe6d99fa712f62f5b6/doc/smi_user_guide.md?plain=1#L637

[qnixsynapse]

I ended up successfully building the package and it shows this output(with root):

+-----------------------------+--------------------------------------------------------------------+
| Device ID                   | 0                                                                  |
+-----------------------------+--------------------------------------------------------------------+
| GPU Utilization (%)         | N/A                                                                |
| EU Array Active (%)         | N/A                                                                |
| EU Array Stall (%)          | N/A                                                                |
| EU Array Idle (%)           | N/A                                                                |
|                             |                                                                    |
| Compute Engine Util (%)     | N/A                                                                |
| Render Engine Util (%)      | Engine 0: 0                                                        |
| Media Engine Util (%)       | N/A                                                                |
| Decoder Engine Util (%)     | Engine 0: 0, Engine 1: 0                                           |
| Encoder Engine Util (%)     | Engine 0: 0, Engine 1: 0                                           |
| Copy Engine Util (%)        | Engine 0: 0                                                        |
| Media EM Engine Util (%)    | Engine 0: 0, Engine 1: 0                                           |
| 3D Engine Util (%)          | N/A                                                                |
+-----------------------------+--------------------------------------------------------------------+
| Reset                       | N/A                                                                |
| Programming Errors          | N/A                                                                |
| Driver Errors               | N/A                                                                |
| Cache Errors Correctable    | N/A                                                                |
| Cache Errors Uncorrectable  | N/A                                                                |
| Mem Errors Correctable      | N/A                                                                |
| Mem Errors Uncorrectable    | N/A                                                                |
+-----------------------------+--------------------------------------------------------------------+
| GPU Power (W)               | 37                                                                 |
| GPU Frequency (MHz)         | 2400                                                               |
| Media Engine Freq (MHz)     | N/A                                                                |
| GPU Core Temperature (C)    | N/A                                                                |
| GPU Memory Temperature (C)  | N/A                                                                |
| GPU Memory Read (kB/s)      | N/A                                                                |
| GPU Memory Write (kB/s)     | N/A                                                                |
| GPU Memory Bandwidth (%)    | N/A                                                                |
| GPU Memory Used (MiB)       | 6804                                                               |
| GPU Memory Util (%)         | 84                                                                 |
| Xe Link Throughput (kB/s)   | N/A                                                                |
+-----------------------------+--------------------------------------------------------------------+

Got the memory usage now but temps are still N/A.

[qnixsynapse]

Adding this here. This may add to the list of todos and fixes. Credit goes to Gemini 1.5 Pro 1 million. :)

• Device-to-device memory copy across GPUs: The current implementation uses a workaround (copying data to host and then back to the other device) for device-to-device memory copies across different GPUs. (This goes along with multi card support)
  [NeoZhangJianyu] Yes, it's known limitation. More efficient implement depends on other library. This function won't impact the performance due to less call times. ✅

• Missing support for certain data types: Some data types, like GGML_TYPE_IQ4_NL, GGML_TYPE_IQ2_S, and GGML_TYPE_IQ4_XS, are not supported in the SYCL backend. Issue: [SYCL] Support newer non linear quantization #5674, PR for GGML_TYPE IQ2_S : [SYCL] iq2_s #6052, A fix (?): IQ1_S: attempt to fix SYCL #6014
  Solved in support/fix more IQ OPs #6521 ✅

• Limited use of SYCL extensions: The code could potentially benefit from utilizing more SYCL extensions for optimization. For example, the ext_oneapi_usm_device_read_only extension could be used for constant memory, and the ext_intel_free_memory extension could provide more accurate memory usage information.
  [AkarshanBiswas] ext_intel_free_memory is used but not sure about ext_oneapi_usm_device_read_only
  [NeoZhangJianyu] Already use them, like ext_intel_free_memory. ✅

• Thread safety: Some parts of the code, like the buffer type map in ggml_backend_sycl_split_buffer_type, might not be thread-safe. This could lead to issues when using multiple models or contexts simultaneously.
  [NeoZhangJianyu] Is there any test case for the issue? I don't test with multiple modes case.
  [AkarshanBiswas] Yes. Test it with continuous batching with slots > 1 and see if the driver is stable.

• Memory management: Implementing a buffer pool for SYCL similar to the CUDA version could improve memory allocation and reuse efficiency.
  [NeoZhangJianyu] The buffer pool is used in SYCL backend. ✅

• Flash attention SYCL implementation after ggml : add Flash Attention #5021

Will update when I find anything else with my limited knowledge.

[airMeng]

#6521

[piDack]

flashattention have any progress?i find huggingface intel backend have impl it https://github.com/intel/intel-extension-for-deepspeed/tree/774453156c4ad6f4b831fb985494254645f5fe1d/intel_extension_for_deepspeed/op_builder/csrc/flash_attn

[slaren]

The SYCL backend should be updated to adopt these changes in ggml-backend:

• Remove the code in ggml.c and implement the offload_op interface: backend : offload large batches to GPU #6083
• Remove the code that copies GGML_BACKEND_TYPE_CPU tensors automatically to VRAM. After adopting the previous change, this will not be necessary, all the tensors received by the SYCL backend will always be allocated in a SYCL buffer.
• Remove all usage of ggml_tensor::backend, as this will be removed in the future. To support split buffer types, use ggml_tensor::buffer to identify the storage type of the tensor instead.

[NeoZhangJianyu]

Yes, we have a plan to fix them by 2 PRs.

[airMeng]

@slaren May I know more context about offload_op? Seems only a ggml_backend_sycl_offload_op needed?

https://github.com/ggerganov/llama.cpp/pull/6217/files

[slaren]

The purpose of offload_op is to offload computation to the GPU when the batch size is large enough, even if the weights are not stored in VRAM and have to be copied. For small batch sizes the cost of copying the weights is higher than the cost of the computation, but for large enough batch sizes, it is often faster to offload the computation to the GPU since it is usually much faster than the CPU.

Previously, this was implemented by hooking into the CPU backend in ggml.c in the function ggml_compute_forward, and doing the computation in a call to ggml_sycl_compute_forward. Now, this is handled in ggml_backend_sched entirely, and the backends only need to implement the offload_op function to choose what operations they wish to run even when the weights are on system memory. The weights are copied to VRAM by ggml_backend_sched, so backends do not need to check if the tensors are in RAM and copy them, they can assume that all the tensors they receive in a call to graph_compute will always be allocated in the backend-specific buffer.

Your implementation is a good first step, but you should also remove the code of SYCL in ggml.c. To test this, you can try prompt processing without fully offloading the model. For example, llama-bench -n 0 -ngl 0. For a backend such as SYCL, it may be good to consider the type of accelerator being used, and possibly disable for slow devices. This will also allow you to remove a lot of the code in the SYCL backend that deals with copying tensors between CPU and GPU, and generally simplify this logic.

You would also need to modify llama.cpp in llama_new_context_with_model to always create an instance of the SYCL backend, even when model->n_gpu_layers == 0.

[airMeng]

@slaren for point 2 and 3, can we refer to #6170 ?

[slaren]

Sure, if you find that useful you can use it. However, there will be further refactoring of the CUDA backend in #6269.

[qnixsynapse]

Just an update here: I did not use llama.cpp for like few days because I was busy. I ran it today to test llama-3 and I found out that it hangs here everytime with every model right here:

......................................................................................
llama_new_context_with_model: n_ctx      = 512
llama_new_context_with_model: n_batch    = 512
llama_new_context_with_model: n_ubatch   = 512
llama_new_context_with_model: freq_base  = 500000.0
llama_new_context_with_model: freq_scale = 1
llama_kv_cache_init:  SYCL_Host KV buffer size =    64.00 MiB
llama_new_context_with_model: KV self size  =   64.00 MiB, K (f16):   32.00 MiB, V (f16):   32.00 MiB
llama_new_context_with_model:  SYCL_Host  output buffer size =     0.98 MiB
llama_new_context_with_model:      SYCL0 compute buffer size =   266.50 MiB
llama_new_context_with_model:  SYCL_Host compute buffer size =    45.01 MiB
llama_new_context_with_model: graph nodes  = 1062
llama_new_context_with_model: graph splits = 66

I am running with --no-mmap .

In the logs I found out:

kernel: Fence expiration time out i915-0000:03:00.0:server[5816]:14c!

Not sure if this is because of an update that I received on Arch Linux or not, not a single model is running with the same binary that it is used to run before.
I have intel-compute-runtime version 23.48.27912.11-2, on lts kernel 6.6.27, intel-oneapi-base-toolkit is 2024.0.0.49564 and level-zero-loader is 1.15.1-1.
Not opening an issue right now because I am not sure if it is a bug in llama.cpp or not.

Update: Not related to llama.cpp, JAX with intel-extension-for-openxla hangs too. (now confirmed)

Update 2: Came across this: intel/compute-runtime#497
Update 3: linux > 6.6.25 is broken (suspicious commit saving it here for bug reporting upstream.)

[qnixsynapse]

Yes. Here is the output of sycl-ls on my PC (I am still trying to debug):

[opencl:acc:0] Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device OpenCL 1.2  [2023.16.10.0.17_160000]
[opencl:cpu:1] Intel(R) OpenCL, 12th Gen Intel(R) Core(TM) i3-12100F OpenCL 3.0 (Build 0) [2023.16.10.0.17_160000]
[opencl:gpu:2] Intel(R) OpenCL Graphics, Intel(R) Arc(TM) A750 Graphics OpenCL 3.0 NEO  [23.48.27912]
[ext_oneapi_level_zero:gpu:0] Intel(R) Level-Zero, Intel(R) Arc(TM) A750 Graphics 1.3 [1.3.27912]

I have reproduced it with llama2-7b-q4 and that's why said earlier than it is not related to llama.cpp. Even JAX hangs. (with intel openxla pjrt plugin) .

[qnixsynapse]

Update: Fixed it!

There is a reason I did not report an issue at first since I was not sure of what was causing the problem. Now I know. And it is the new linux-firmware version: 20240409.1addd7dc-1 linux version > 6.6.25 .

I downgraded to version 6.6.25 and now it works.

Thinking of switching to debian on this system, I really hate updates breaking my important things around.

[airMeng]

can you help to report the issue to https://github.com/oneapi-src/level-zero? much thanks for your help!

[qnixsynapse]

No problem!
Yes, I will open an issue tomorrow.

[qnixsynapse]

Update: The regression happened on i915 kernel driver and potential commits(according to Matt):

a7ff84a6fe5a drm/i915/gt: Enable only one CCS for compute workload
726ff623869d drm/i915/gt: Do not generate the command streamer for all the CCS
c1f7ce2a11a9 drm/i915/gt: Disable HW load balancing for CCS

Quoting him, "That workaround requires that we disable load balancing on the compute engines, assign all underlying hardware compute resources to a single engine, and then hide all except for one of the engines from userspace (which is also why the "physical engines" count goes from 4 down to 1 as you noted). It's a pretty complicated and invasive hardware workaround, so it has the highest likelihood of accidentally introducing issues."

Fix is 4cfca03f7641 ("drm/i915/gt: Automate CCS Mode setting during engine resets") which will hopefully land on stable kernels by next week.

Adding these here so that people who are experiencing similar issues might find these helpful.

[airMeng]

Latest SYCL is broken due to #7640 (comment), I am looking into it and hope to fix it soon.

cc @NeoZhangJianyu @AidanBeltonS

[airMeng]

fixed in #7710

[mistrjirka]

I am not sure if it is expected or not but attempting to use Intel GPU ( Intel® Iris® Xe Graphics ) on Intel® Core™ i7-1165G7 results in much slower performance then on the CPU alone. I tried just the example prompt:

ZES_ENABLE_SYSMAN=1 ./build/bin/llama-cli -m models/Lexi-Llama-3-8B-Uncensored_Q4_K_M.gguf -p "Building a website can be done in 10 simple steps:" -n 400 -e -ngl 33 -sm none -mg 0
I get:

So around 1.6t/s for prompt eval and 3.55t/s for generation.
The GPU is used according to intel_gpu_top:

When I try it on CPU I get much better performance:
command: ZES_ENABLE_SYSMAN=1 ./build/bin/llama-cli -m models/Lexi-Llama-3-8B-Uncensored_Q4_K_M.gguf -p "Building a website can be done in 10 simple steps:" -n 400 -e -ngl 0 -sm none -mg 0

It also seems that the output from the GPU is garbled and generally wrong.

[airMeng]

hi @mistrjirka we have observed the extremely slow prompt evel issue and we are working on it. As for prompt generation, as you might know, iGPU and CPU shares the same DDR memory so the performance shall be on the same level. we are working on it too but expect not too much gains.

cc @luoyu-intel

[mistrjirka]

hi @mistrjirka we have observed the extremely slow prompt evel issue and we are working on it. As for prompt generation, as you might know, iGPU and CPU shares the same DDR memory so the performance shall be on the same level. we are working on it too but expect not too much gains.

cc @luoyu-intel

Thanks for the info and thank you for working on this.
I am honestly more concerned about the weird repeating output
Tried the prompt multiple times and the output is generally just nonsense. Sometimes it print random characters. It shows up even when using less gpu layers. That could point to some performance non related problem. I observed that it holds on Intel xe graphics and on Intel 620m. I am not sure if someone can replicate my problems on dedicated Card.

[airMeng]

we are aware of that too, seems the recent update of rope and mul_mat_id breaks the accuracy. we are working on it, sorry for any inconvenience.

[qnixsynapse]

Hi. I have some questions.

Is this intentional?

We can simply call the functions with the required parameters which should be better right?

[qnixsynapse]

I think we should simplify this so that function for each operation can be independent in choosing its parameters.

[qnixsynapse]

Thoughts?
Also, CC @Rbiessy and maybe @abhilash1910

[Rbiessy]

On the first question about refactoring ggml_sycl_compute_forward it seems that func is used to be able to call ggml_sycl_set_peer_access in some cases before calling func. Assuming we want to keep this behavior I would suggest to keep func but to change the type of ggml_sycl_func_t so that it does not contain src0 and src1 arguments. They can all be deduced from the last ggml_tensor object.

On the second question about using mmq, we could help to review this PR. It has some large implications:

• We would need to verify the correctness of such a pass. Unfortunately running the tests is not always enough, internally we also run the model and verify that the output is not gibberish.
• We would need to verify the impact on performance for some models and some Nvidia and Intel devices. The PR would need to show evidence that it has a overall a positive impact. We should be able to help with that.

[qnixsynapse]

Thank you for sharing your insights. My original plan of refactoring compute_forward function was shelved because I decided to use GGML_UNUSED() macro for unused variables.

My current plan is to make this backend migrate to tensor->buffer and write optimized mmq kernels for Intel to speed up PP. But I am not an expert in writing SYCL kernels so I not having enough confidence to do so.

The screenshot shows that Q4_K, Q5_K and Q6_K kernels are broken, hence I have commented them out because we did this in the last PR which got merged:
https://github.com/ggerganov/llama.cpp/blob/64ae0655114f84f11a724bc6878c6f8f4a55560b/ggml/src/ggml-sycl/mmq.cpp#L816-L817

Also, no support for IQ_XX data types because these kernels were never updated.

Edit: According to this, Xe-HPG architecture has a register size of 32 KB. RDNA architecture, whose tiling settings I found in SYCL's mmq kernels has 128KB of register size. I am suspecting that register overflows may have been causing accuracy issues.

https://github.com/ggerganov/llama.cpp/blob/9f35e44592a7646a5803620eb6a3f0ed5ac90553/ggml/src/ggml-sycl/mmq.cpp#L1785-L1803

[qnixsynapse]

On the first question about refactoring ggml_sycl_compute_forward it seems that func is used to be able to call ggml_sycl_set_peer_access in some cases before calling func. Assuming we want to keep this behavior I would suggest to keep func but to change the type of ggml_sycl_func_t so that it does not contain src0 and src1 arguments. They can all be deduced from the last ggml_tensor object.

Actually, it checks if the src0 tensor is of split buffer type and then it calls ggml_sycl_set_peer_access. I initially thought to leave it as it is but it seems I have no choice but to refactor it.

The current code assumes that all the operation needs two src tensors (i.e src0 and src1) and because of that I find it difficult to add new OPs. Not all OPs need src0 and src1 tensors. Some need 4 src tensors such as flash attention (src0 as Q, src1 as K, src2 as V and src3 as mask), along with extra params such as bias value(for alibi) and logits softcap(Gemma models). Everything is there in node/dst tensor. The best is to check if src0 is present and is a split buffer and run ggml_sycl_set_peer_access before the switch case statement IMO.

[Mar2ck]

I'd really like to see SYCL support added to the RPC server.

[qnixsynapse]

SIGBUS {si_signo=SIGBUS, si_code=BUS_ADRERR, si_addr=0x74f11d03e020} ---

This one is helpful but incomplete. You have to rebuild the project with debug symbols by either adding -g to C++ flags / replacing --config Release with --config Debug and run it in a program called gdb.

I know it's a bit technical. I hope situation will improve once my refactoring work is done and merged which will correctly print the backend debug logs.

[rtlinux]

Thanks for the tips! Something like this...?

load_tensors: layer  23 assigned to device SYCL0
load_tensors: layer  24 assigned to device SYCL0
load_tensors: tensor 'token_embd.weight' (q8_0) (and 0 others) cannot be used with preferred buffer type CPU_AARCH64, using CPU instead
load_tensors: offloading 24 repeating layers to GPU
load_tensors: offloading output layer to GPU
load_tensors: offloaded 25/25 layers to GPU
load_tensors: RPC[192.168.1.178:50053] model buffer size =   129.39 MiB
load_tensors:   CPU_Mapped model buffer size =   137.94 MiB
load_tensors:        SYCL0 model buffer size =   244.35 MiB

Thread 1 "llama-cli" received signal SIGBUS, Bus error.
__memcpy_avx_unaligned_erms () at ../sysdeps/x86_64/multiarch/memmove-vec-unaligned-erms.S:540
warning: 540    ../sysdeps/x86_64/multiarch/memmove-vec-unaligned-erms.S: No such file or directory
(gdb)

backtrace:

[qnixsynapse]

0x00007ffff7a9da7f in ggml_backend_sycl_buffer_set_tensor(ggml_backend_buffer*, ggml_tensor*, void const*, unsigned 
long, unsigned long)

aligns with

mmap(NULL, 256245760, PROT_READ|PROT_WRITE, MAP_SHARED, 4, 0x156304000) = 0x74f11d03e000
--- SIGBUS {si_signo=SIGBUS, si_code=BUS_ADRERR, si_addr=0x74f11d03e020} ---

Does this happen even on llama-cli or llama-server too?

Also, proper way to get debug logs is to first set pagination off and bt full in gdb.

edit: Also check ReBar settings in your bios.

[rtlinux]

ReBar setting was indeed the culprit, it's all working now after turning it on in the bios. Thanks again for your help!

[qnixsynapse]

No problem!. ;)

[NeoZhangJianyu]

I think the internal refactor should support the outside features to user.
I suggest to move the focus to optimize the kernel performance on Intel GPU.

[Rbiessy]

I am planning to work on renaming the occurrences of oneMKL Interface to oneMath following the recent renaming.

Currently oneMath is used for the Nvidia and AMD devices when SYCL is enable while the Intel devices directly use the Intel oneMKL product, see this FAQ for the difference between the two.. It would be valuable to also use oneMath for Intel devices (which will directly call Intel oneMKL) to use a common interface. I believe this will not impact performance if we use the compile time dispatcher.
One issue is that oneMath is not shipped with pre-built packages which can make building llama.cpp harder. The best solution to mitigate this is to provide 2 ways to use oneMath:

1. Use the optional installation provided by users with -DoneMath_DIR=/path/to/oneMath/install
2. Use CMake FetchContent to automatically pull and build oneMath if SYCL is enabled and oneMath cannot be found. This should not add more requirements for the user. Building oneMath BLAS domain with Intel oneMKL backends only takes 3s on my desktop.

Does anyone have a concern with this approach? In particular curious for @airMeng and @NeoZhangJianyu opinions.

[qnixsynapse]

Personally, I do not have problem with this. But we need to setup a proper CI(if not Intel, then Nvidia) and write some runtime tests here imo.

[Rbiessy]

More CI is probably good but I'm not sure how it relates to this suggestion.
I don't know if the public CI has the HW available to run some runtime tests. It's not a priority for us today but a nice to have.

[qnixsynapse]

Last time I checked, Intel CI had some issues here. oneMath reminded me of the idea that we can use already available HW here for CI. Nevermind.

[Rbiessy]

I have implemented the suggested changes in #12192

[qnixsynapse]

@Rbiessy @Alcpz @NeoZhangJianyu @airMeng , I am noticing a memory leak with each token output only when building with SYCL backend. I am unable to use vtune for the following reason on an up-to-date Fedora:

Building sep driver ...
In file included from /usr/src/kernels/6.13.5-200.fc41.x86_64/include/linux/module.h:22,
                 from lwpmudrv.c:37:
lwpmudrv.c:106:18: error: expected ‘,’ or ‘;’ before ‘INTEL_PMT’
  106 | MODULE_IMPORT_NS(INTEL_PMT);
      |                  ^~~~~~~~~
/usr/src/kernels/6.13.5-200.fc41.x86_64/include/linux/moduleparam.h:26:61: note: in definition of macro ‘__MODULE_INFO’
   26 |                 = __MODULE_INFO_PREFIX __stringify(tag) "=" info
      |                                                             ^~~~
/usr/src/kernels/6.13.5-200.fc41.x86_64/include/linux/module.h:299:33: note: in expansion of macro ‘MODULE_INFO’
  299 | #define MODULE_IMPORT_NS(ns)    MODULE_INFO(import_ns, ns)
      |                                 ^~~~~~~~~~~
lwpmudrv.c:106:1: note: in expansion of macro ‘MODULE_IMPORT_NS’
  106 | MODULE_IMPORT_NS(INTEL_PMT);
      | ^~~~~~~~~~~~~~~~
lwpmudrv.c:107:18: error: expected ‘,’ or ‘;’ before ‘INTEL_PMT_TELEMETRY’
  107 | MODULE_IMPORT_NS(INTEL_PMT_TELEMETRY);
      |                  ^~~~~~~~~~~~~~~~~~~
/usr/src/kernels/6.13.5-200.fc41.x86_64/include/linux/moduleparam.h:26:61: note: in definition of macro ‘__MODULE_INFO’
   26 |                 = __MODULE_INFO_PREFIX __stringify(tag) "=" info
      |                                                             ^~~~
/usr/src/kernels/6.13.5-200.fc41.x86_64/include/linux/module.h:299:33: note: in expansion of macro ‘MODULE_INFO’
  299 | #define MODULE_IMPORT_NS(ns)    MODULE_INFO(import_ns, ns)
      |                                 ^~~~~~~~~~~
lwpmudrv.c:107:1: note: in expansion of macro ‘MODULE_IMPORT_NS’
  107 | MODULE_IMPORT_NS(INTEL_PMT_TELEMETRY);

Memory grows 2GB per requests to the server. I tested with other backends(Vulkan) and it doesn't happen in there.

This started happening recently, I am not sure when. Can you all please help me in identifying it? If I run with valgrind, the entire thing crashes.

[qnixsynapse]

Found the culprit... It's this: 08d5986

Me and @Alcpz was thinking of reverting it because it somehow breaks RPC server on Q4_0. But I left it because I thought it is a good optimisation implementation. But it seems now that it breaks more than just RPC on q4_0

The model I was testing on was Q5_0 and others and not Q4_0 which triggers the leak.

[Alcpz]

Yes. I am with @qnixsynapse. I felt reluctant of reverting because we get such good performance. Can we try to find where the leak happens and fix that?

[qnixsynapse]

@NeoZhangJianyu Since you are the author, so let me ask you.

Why are we adding element to the vector tensor_extras but never freeing it, even on a non Q4_0 model?

llama.cpp/ggml/src/ggml-sycl/ggml-sycl.cpp

Lines 337 to 339 in e0dbec0

	ggml_tensor_extra_gpu * extra = new ggml_tensor_extra_gpu{};
	tensor->extra = extra;
	ctx->tensor_extras.push_back(extra); //used to release it when destroy ctx.

[qnixsynapse]

Opened #12366

[NeoZhangJianyu]

Oh, it's my fault! :)
The extra is used to record the tensor state for reorder.
It will be released with the context release.
I think it may be the root cause.

Thank you to fix by #12366.

Need to free it in reset function as comment.

[jeffzhou2000]

the technical atmosphere of this Intel sycl is very good and I really envy it. Intel is a really great IT company. I'd like to spend some time learning knowledge about Intel sycl in the future.

[NeoZhangJianyu]

Welcome to join us! :)

[jeffzhou2000]

thanks so much for your warmly welcome and I will start as a beginner at the end of this month and setting up the dev envs on my Intel x86 desktop SoC equipped Linux machine.

[jeffzhou2000]

I'm sorry that I couldn't start the learning&dev activity on ggml-sycl since the end of this month because it seems there is another tech approach of ggml-hexagon(llama.cpp backend for Qualcomm's Hexagon NPU), I'd like to try it regardless of success of failure.

I'll do my commitment on ggml-sycl because(1)I learnt too much from ggml-sycl on 04/2024 (2)Intel is a really great IT company.

[NeoZhangJianyu]

OK! It's good!
llama.cpp is a great platform to integrate many technology for AI.
It's easy to help people learn them. :)

[qnixsynapse]

Just for information:
While trying to setup SYCL CI on llama.cpp repo, I found out that commit 08d5986 broke Q4_0 Pythia model which is used in SYCL CI. Using GGML_SYCL_DISABLE_OPT=1 fixes it.

[NeoZhangJianyu]

Could you share the CI error log?
I only see 1 CPY case is fault in my local test.

[NeoZhangJianyu]

When I run the latest CI cases, all are passed.
Could you share your fault cases?

[qnixsynapse]

When I run the latest CI cases, all are passed.
Could you share your fault cases?

Have you tested the output generation of this model?

[qnixsynapse]

Full context is here (reproducible on a single GPU as well) #12546 (comment)

[NeoZhangJianyu]

let me check.

[ky438]

@airMeng - feature request: when will llama.cpp support the DPAS/DPASW instructions to perform efficient matrix multiplication on Intel GPUs? Perhaps by using the SYCL joint_matrix extension?

[ShanoToni]

Questions on SYCL MUL_MAT Reorder Optimizations

Hello, I would like to request some clarification of the Q4_0 reorder optimisation.

• When creating a temporary copy of the tensor for reordering, the temporary memory is allocated as shared memory. Since this is a device to device copy wouldn’t it be more performant to do malloc_device?

  llama.cpp/ggml/src/ggml-sycl/ggml-sycl.cpp

  Line 3566 in 75afa0a

  auto tmp_buf = sycl::malloc_shared<char>(size, *stream);

• The backend context sets a bool that it has optimised a graph the first time optimize_graph_once is called. Any following calls to compute a graph with the same context, for example in the test_backend_ops does not optimise the graph, meaning that only the first test case is passing. Running the tests with GGML_SYCL_DISABLE_OPT=0 and commenting out this line causes failures in Mul_mat test cases

  llama.cpp/ggml/src/ggml-sycl/ggml-sycl.cpp

  Line 3619 in 75afa0a

  ctx->optimized_graph = true;

• Is the addition of the ggml_tensor_extra_gpu required? AFAIK this is a discouraged approach to optimisations. An alternative would be to to create a heuristic function similar to

bool reorder_tensor_case(tensor * tensor_to_reorder, tensor * other, tensor * result) {
    if (!g_ggml_sycl_disable_optimize && result->op == GGML_OP_MUL_MAT && tensor_to_reorder->type == GGML_TYPE_Q4_0 &&
            other->ne[2] == 1 && other->ne[3] == 1 && tensor_to_reorder->op == GGML_OP_NONE) {
            return true;
    }    
    return false;
}

• This can be used
  • Here:

    llama.cpp/ggml/src/ggml-sycl/ggml-sycl.cpp

    Line 3605 in 75afa0a

    if (dst->op == GGML_OP_MUL_MAT && src0->type == GGML_TYPE_Q4_0 &&

  • And here:

    llama.cpp/ggml/src/ggml-sycl/dmmv.cpp

    Line 1106 in 75afa0a

    if ((ggml_tensor_extra_gpu*)dst->src[0]->extra &&

This would remove the need to modify the reordered tensor kernel with the extra

• Lastly the reorder seems to not take into account the scenario where the Q4_0 tensor is the output of an Operation. This seems related to the failures mentioned above, In one of the failing cases the mul_mat is configured with a permute. While this is a no-op, it does change the mul_mat kernel being run from the dequantize_reorder to another configuration, which does not support/expect the reordered tensor format resulting in the incorrect results
• Example of failing config

test_mul_mat(GGML_TYPE_Q4_0, GGML_TYPE_F32, 16, 2, 256, {1,  1}, {1, 1}, {0, 2, 1, 3})

To resolve this in the proposed reorder_tensor_case func an additional check was added

 tensor_to_reorder->op == GGML_OP_NONE

This should ensure that the tensor is not the result of another permutation.

Please let me know your thoughts on these topics, if any of these make sense I would be happy to raise a PR with the proposed changes.

[NeoZhangJianyu]

• When creating a temporary copy of the tensor for reordering, the temporary memory is allocated as shared memory. Since this is a device to device copy wouldn’t it be more performant to do malloc_device?

A:malloc_device and malloc_share have same performance when copy memory in same device.
In this case, there is no device to device copy.

• The backend context sets a bool that it has optimised a graph the first time optimize_graph_once is called. Any following calls to compute a graph with the same context, for example in the test_backend_ops does not optimise the graph, meaning that only the first test case is passing. Running the tests with GGML_SYCL_DISABLE_OPT=0 and commenting out this line causes failures in Mul_mat test cases

A: No. test_backend_ops trigger the optimize_graph_once too.
All mul_mat cases should be passed when enable reorder optimize. At least it's passed with commit for reorder optimize feature.

• Is the addition of the ggml_tensor_extra_gpu required? AFAIK this is a discouraged approach to optimisations. An alternative would be to to create a heuristic function similar to

A: Yes, it's required.
The function you draft can't make sure the reorder is used in the tensor for all cases/LLMs.
Use extra to record by a flag is simple and accuracy method.

• test_mul_mat(GGML_TYPE_Q4_0, GGML_TYPE_F32, 16, 2, 256, {1, 1}, {1, 1}, {0, 2, 1, 3})

A:
This part will be refactored by me. It's WIP.

I just test, all UT cases are passed when enable this feature.

Not all Ops with Q4_0 are optimized by reorder feature.
In current code, I only implement for a part of cases of Q4_0.
We will extend it to all Q4_0 and Q4,Q5,..Q8.

[ShanoToni]

Thank you @NeoZhangJianyu for taking the time and answering the questions. I wanted to expand on a few of the points.

1. I wanted to ask, Is this optimization unique for SYCL or are you aware of any other backend using a similar reorder? If so would you please point me to it.
2. Regarding the UT passing. I do not believe that the UT use the reordered tensor. If you print the context attempting to optimize the graph Here

   llama.cpp/ggml/src/ggml-sycl/ggml-sycl.cpp

   Line 3628 in 75afa0a

   if (!g_ggml_sycl_disable_optimize) optimize_graph_once(cgraph, sycl_ctx);

   you get the same pointer, meaning that all tests are run with the same context. The early return here

   llama.cpp/ggml/src/ggml-sycl/ggml-sycl.cpp

   Line 3616 in 75afa0a

   if (ctx->optimized_graph) {

   and setting the optimized_graph = true would mean that the context would only attempt to optimize the first graph and NOT any following graphs run with the same context. In order to prove that this is the case I modified the code here

   llama.cpp/ggml/src/ggml-sycl/ggml-sycl.cpp

   Line 3616 in 75afa0a

   if (ctx->optimized_graph) {

   to

    if (ctx->optimized_graph) {
        for (int i = 0; i < cgraph->n_nodes; i++) {
            ggml_tensor *src0 = cgraph->nodes[i]->src[0];
            ggml_tensor *src1 = cgraph->nodes[i]->src[1];

            if (cgraph->nodes[i]->op == GGML_OP_MUL_MAT && src0->type == GGML_TYPE_Q4_0 &&
                src1->ne[2]==1 && src1->ne[3]==1) {
                    printf("\\nContext: %p\\n ", ctx);
                    printf("Current graph node should be optimised but is not due to ctx->optimized_graph = true\\n");
                }
        }
       return;
    }

The added if checks for the same conditions as here

llama.cpp/ggml/src/ggml-sycl/ggml-sycl.cpp

Line 3605 in 75afa0a

if (dst->op == GGML_OP_MUL_MAT && src0->type == GGML_TYPE_Q4_0 &&

and reports if the graph node would be optimized if the context.optimized_graph was not set to true . Running this trough the MUL_MAT tests we get the output of:

MUL_MAT(type_a=f32,type_b=f32,m=16,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f32,type_b=f32,m=16,n=2,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f32,type_b=f32,m=16,n=3,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f32,type_b=f32,m=16,n=4,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f32,type_b=f32,m=16,n=5,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f32,type_b=f32,m=16,n=6,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f32,type_b=f32,m=16,n=7,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f32,type_b=f32,m=16,n=8,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f32,type_b=f32,m=16,n=9,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f16,type_b=f32,m=16,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f16,type_b=f32,m=16,n=2,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f16,type_b=f32,m=16,n=3,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f16,type_b=f32,m=16,n=4,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f16,type_b=f32,m=16,n=5,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f16,type_b=f32,m=16,n=6,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f16,type_b=f32,m=16,n=7,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f16,type_b=f32,m=16,n=8,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=f16,type_b=f32,m=16,n=9,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): OK
  MUL_MAT(type_a=bf16,type_b=f32,m=16,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): not supported [SYCL0]
  MUL_MAT(type_a=bf16,type_b=f32,m=16,n=2,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): not supported [SYCL0]
  MUL_MAT(type_a=bf16,type_b=f32,m=16,n=3,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): not supported [SYCL0]
  MUL_MAT(type_a=bf16,type_b=f32,m=16,n=4,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): not supported [SYCL0]
  MUL_MAT(type_a=bf16,type_b=f32,m=16,n=5,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): not supported [SYCL0]
  MUL_MAT(type_a=bf16,type_b=f32,m=16,n=6,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): not supported [SYCL0]
  MUL_MAT(type_a=bf16,type_b=f32,m=16,n=7,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): not supported [SYCL0]
  MUL_MAT(type_a=bf16,type_b=f32,m=16,n=8,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): not supported [SYCL0]
  MUL_MAT(type_a=bf16,type_b=f32,m=16,n=9,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): not supported [SYCL0]
  MUL_MAT(type_a=q4_0,type_b=f32,m=16,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): Context: 0x35797530
 Current graph node should be optimised but is not due to ctx->optimized_graph = true
OK
  MUL_MAT(type_a=q4_0,type_b=f32,m=16,n=2,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): Context: 0x35797530
 Current graph node should be optimised but is not due to ctx->optimized_graph = true
OK
  MUL_MAT(type_a=q4_0,type_b=f32,m=16,n=3,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): Context: 0x35797530
 Current graph node should be optimised but is not due to ctx->optimized_graph = true
OK
  MUL_MAT(type_a=q4_0,type_b=f32,m=16,n=4,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): Context: 0x35797530
 Current graph node should be optimised but is not due to ctx->optimized_graph = true
OK
  MUL_MAT(type_a=q4_0,type_b=f32,m=16,n=5,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): Context: 0x35797530
 Current graph node should be optimised but is not due to ctx->optimized_graph = true
OK
  MUL_MAT(type_a=q4_0,type_b=f32,m=16,n=6,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): Context: 0x35797530
 Current graph node should be optimised but is not due to ctx->optimized_graph = true
OK
  MUL_MAT(type_a=q4_0,type_b=f32,m=16,n=7,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): Context: 0x35797530
 Current graph node should be optimised but is not due to ctx->optimized_graph = true
OK
  MUL_MAT(type_a=q4_0,type_b=f32,m=16,n=8,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): Context: 0x35797530
 Current graph node should be optimised but is not due to ctx->optimized_graph = true
OK
  MUL_MAT(type_a=q4_0,type_b=f32,m=16,n=9,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0): Context: 0x35797530
 Current graph node should be optimised but is not due to ctx->optimized_graph = true
OK
. . .

Meaning that any of the following test cases are NOT reordered, since the context attempt to optimize only the first graph for the mul_mat (type_a=f32,type_b=f32,m=16,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],v=0) which is not a valid reorder case( type_a is required to be q4_0)

3. Another problem I have noticed investigating this is a difference in selecting when to do the reoder optimization

   llama.cpp/ggml/src/ggml-sycl/ggml-sycl.cpp

   Line 3605 in 75afa0a

   if (dst->op == GGML_OP_MUL_MAT && src0->type == GGML_TYPE_Q4_0 &&

   and the condition for running the ggml_sycl_op_dequantize_mul_mat_vec https://github.com/ggml-org/llama.cpp/blob/75afa0ae31f0a51aaadcc5ff146eb7a32a7f9088/ggml/src/ggml-sycl/ggml-sycl.cpp#L2889.The first if checks ne[2]==1 and ne[3]==1 while the use_dequantize_mul_mat_vec requires ne[1]==1 . This would mean that in a case where the ne = {K, N, 1,1 } and N > 1 the tensor will be reordered BUT the mul mat called would be ggml_sycl_op_mul_mat_vec_q or ggml_sycl_op_mul_mat_q instead of ggml_sycl_op_dequantize_mul_mat_vec which is currently the only one which supports the reordered tensor.

We are planning to work on submitting PRs that fixes both of the issues mentioned because it is starting to become an issue for us and it could help the discussion.

[qnixsynapse]

A: Yes, it's required.
The function you draft can't make sure the reorder is used in the tensor for all cases/LLMs.
Use extra to record by a flag is simple and accuracy method.

I strictly discourage setting tensor extras in init_tensor. IMHO, we should find an alternative solution.

Edit:
Can't these things be checked though a backend buffer ctx interface?

llama.cpp/ggml/src/ggml-sycl/common.hpp

Lines 278 to 309 in b43d89e

	struct ggml_tensor_extra_gpu {
	void* data_device[GGML_SYCL_MAX_DEVICES]; // 1 pointer for each device for split
	// tensors
	dpct::event_ptr events[GGML_SYCL_MAX_DEVICES]
	[GGML_SYCL_MAX_STREAMS]; // events for synchronizing multiple GPUs
	optimize_feature optimized_feature;
	};
	void release_extra_gpu(ggml_tensor_extra_gpu * extra, std::vector<queue_ptr> streams={});
	inline optimize_feature check_gpu_optimize_feature(syclex::architecture &arch) {
	optimize_feature opt;
	opt.reorder =
	(arch == syclex::architecture::intel_gpu_dg1 ||
	arch == syclex::architecture::intel_gpu_acm_g10 ||
	arch == syclex::architecture::intel_gpu_acm_g11 ||
	arch == syclex::architecture::intel_gpu_acm_g12 ||
	arch == syclex::architecture::intel_gpu_pvc ||
	arch == syclex::architecture::intel_gpu_pvc_vg ||
	arch == syclex::architecture::intel_gpu_mtl_u ||
	arch == syclex::architecture::intel_gpu_mtl_s ||
	arch == syclex::architecture::intel_gpu_mtl_h ||
	arch == syclex::architecture::intel_gpu_arl_u ||
	arch == syclex::architecture::intel_gpu_arl_s ||
	arch == syclex::architecture::intel_gpu_arl_h ||
	arch == syclex::architecture::intel_gpu_bmg_g21 ||
	arch == syclex::architecture::intel_gpu_lnl_m
	);
	return opt;
	}

[slaren]

The CPU backend implements weight reordering by adding a different buffer type, which is exposed via the ggml_backend_dev_get_extra_bufts optional function. llama.cpp will automatically try to use the buffer types returned by this function when loading the weights (*). Generally this would be the recommended way to implement this, modifying the tensors in the call to graph_compute is likely to result in all kinds of problems.

(*) Actually this is skipped for GPU backends at the moment, however there is no reason to not enable it for GPU backends as well.

[NeoZhangJianyu]

The CPU backend implements weight reordering by adding a different buffer type, which is exposed via the ggml_backend_dev_get_extra_bufts optional function. llama.cpp will automatically try to use the buffer types returned by this function when loading the weights (*). Generally this would be the recommended way to implement this, modifying the tensors in the call to graph_compute is likely to result in all kinds of problems.

(*) Actually this is skipped for GPU backends at the moment, however there is no reason to not enable it for GPU backends as well.

Yes, it's good suggestion.
I will consider it.

[NeoZhangJianyu]

@ShanoToni
Your questions are too long, I try to answer them one by one.
Please don't list too many questions in one post. :)

• I wanted to ask, Is this optimization unique for SYCL or are you aware of any other backend using a similar reorder? If so would you please point me to it.

A: Currently, it's unique. I don't try in other backend.
It's only approved work well in SYCL backend on Intel GPUs (exclude old iGPU in Core CPU).
Only perfect implementation can approve it work well in other backend in the hardware.
It depends on the hardware features and program language, compiler.

• Regarding the UT passing. I do not believe that the UT use the reordered tensor.

A: I feel it's wrong. But I need to check the code later. I will answer you after checking.

• Another problem I have noticed investigating this is a difference in selecting when to do the reoder optimization

A: Yes, your concern is right.
It will lead to the result error of some LLMs. That's why reorder feature is changed to disable as default.
I'm fixing this issue.
If your work is in beginning, I suggest pending it.

[ShanoToni]

Thanks again for taking the time and answering, and yes in future I'll make a point of separating my points in separate comments so its easier :)

[NeoZhangJianyu]

• Regarding the UT passing. I do not believe that the UT use the reordered tensor.

I don't want to check it, since I'm fixing the known issue of reorder feature.
I will make sure the UT cases will cover the reorder any way.

Thank you!

[qnixsynapse]

Regarding the UT passing. I do not believe that the UT use the reordered tensor.

Probably I disabled it because of accuracy issues. #12560

[NeoZhangJianyu]

I create a PR: #13003 to fix all known issues.
All you concern will be matched. Including issue in #12560.

[ShanoToni]

Hello, we wanted to share our investigation regarding concerns about the reorder optimization using the extra field and potential problems with RPC and getting it more inline with the way CPU backends would repack their tensors.

TL;DR We looked into removing the extra and be more inline with CPU approach, it does not seem doable in a way that would not impede performance benefits of reorder and might introduce different problems for the RPC case.

First major difference is that CPU backends seem to make the change of buffer types based on the defined architecture of the CPU. This is done regardless of tensor shape, this is not the case for the SYCL backend. Here:

llama.cpp/ggml/src/ggml-cpu/ggml-cpu.cpp

Line 36 in 514c456

std::vector<ggml_backend_buffer_type_t>& ggml_backend_cpu_get_extra_buffers_type() {

Additionally it seems that the CPU backend also leverages the extra field to annotate their tensors with specific traits for repacking. Here:

llama.cpp/ggml/src/ggml-cpu/ggml-cpu-aarch64.cpp

Line 6407 in 514c456

ggml::cpu::tensor_traits * get_tensor_traits(const struct ggml_tensor * op) override {

At tensor_init this extra is set which determines which repack type should be used:

llama.cpp/ggml/src/ggml-cpu/ggml-cpu-aarch64.cpp

Line 6323 in 514c456

static enum ggml_status ggml_backend_cpu_aarch64_buffer_init_tensor(ggml_backend_buffer_t buffer, struct ggml_tensor * tensor) {

At tensor_set the incoming data is repacked based on the type set in init

llama.cpp/ggml/src/ggml-cpu/ggml-cpu-aarch64.cpp

Line 6336 in 514c456

auto OK = tensor_traits->repack(tensor, data, size);

MulMat is run with this optmized repack shape of the tensor quants. This is similar to what the reorder optimization for SYCL attempt to do with some differences.

• First differences is the extra buffer types are kept inside the context of the ggml_backend_buffer_t for SYCL this context is already used to contain the device index, name and stream(queue)
• Secondly the we have investigated a way of potentially inlining the approach with the CPU implementation, which we we concluded would not improve the current approach
  • We attempted to move the reorder extra inside the tensor_init, if conditions are met for the mul_mat dst tensor (there are shape and quantization constraints for the reorder to be supported) the src tensor's extra is set to reorder and at tensor_set if the extra is set the incoming data is reordered and stored into the tensor in that fashion. This however does not solve the problem of the extra field being used.
  • In order to do this, since we have conditions of the mul_mat shape to reorder, we would need to change the buffer interface of the tensor we would like to reorder, and modify the get/set/cpy tensor functions to reorder the data before reading it in/out of the tensor. This approach would work HOWEVER this would break for RPC.
  • Since the mul_mat checks the extra field when launching the kernel do determine if it should launch in a reordered fashion or not not having the extra would mean we would lose a way of determining which mul_mat kernel needs to be run.

    llama.cpp/ggml/src/ggml-sycl/dmmv.cpp

    Line 1106 in 514c456

    if ((ggml_tensor_extra_gpu*)dst->src[0]->extra &&

• The current reorder optimization approach does not have this caveat. Additionally in the RPC case since the extra field is allocated at tensor_init time and has null checks at every use, the reorder optimization would simply not occur in the RPC case and the mulmat would simply work in default case.

[slaren]

I think this should be enough to address these issues:

• You can use the supports_op function with any logic necessary to decide if the tensor should be repacked, shape, type, or anything else
• You can use the tensor buffer type to determine which kernel to run, ie. in the same way it is done with ggml_backend_buffer_is_sycl_split

Note that the RPC does not use the extra buffer types, so this wouldn't allow to use tensor reordering with RPC servers. However, it would still allow it to work without reordering, and in the future the RPC backend may add support for forwarding the extra buffer type.

[NeoZhangJianyu]

mul_mat for Q4_0 of SYCL backend has several kernel functions for different shape of src0, src1.
Currently, reorder feature of SYCL backend is not finished. some of above kernel functions don't support reorder.
So when reorder the src0, some cases will be skipped to reorder in a LLM inference.

So suggest RPC disable the reorder in current stage, by g_ggml_sycl_disable_optimize=1 .
After all Q4_0 kernel functions of mul_mat of SYCL support reorder, RPC could enable reorder for all Q4_0 mul_mat.

[Rbiessy]

@NeoZhangJianyu following your comment in #13724 (comment), you mentioned the reorder optimization is "not supported" on 10th gen iGPU, do you know why that is? Do you maybe mean that this optimization makes the performance worse on older devices?
I have some concerns with the current checks in check_gpu_optimize_feature as newer devices won't enable the reorder opt by default which is an issue. I was thinking of changing this so the reorder opt is enabled by default and only disabled for devices we know it performs worse.

[NeoZhangJianyu]

I have tested it.
It can't bring performance improvement.
Disable to support it will avoid reorder the weights for additional cost.

The detection devices doesn't cover all existed Intel GPUs.
Set it as default for unknow GPU, will enable it in older GPUs too.

I suggest using white list for this feature enable.
It includes all GPUs verified to get plus.

[kwaa]

Just noticed that the action to publish a SYCL docker image has been failing for two weeks now.

https://github.com/ggml-org/llama.cpp/actions/runs/15458334203/job/43514658105

[Rbiessy]

Thanks for reporting this, it is now fixed with #14144

[HumerousGorgon]

Does anybody have an idea why enabling flash attention with llama-server completely nukes performance? If I'm understanding it correctly, it's because the SYCL implementation doesn't have the right kernels implemented to support it, or is it due to something else?

[qnixsynapse]

Because flash attention is not yet supported so if you enable it in SYCL, ggml offloads the layers to CPU to do prompt processing and CPU is slow.

[slaren]

The SYCL backend seems to fail occasionally the test test-thread-safety added in #14035.
https://github.com/ggml-org/ci/tree/results/llama.cpp/5b/10edfe34390d546589a6e752e341669c76aa64/ggml-6-x86-sycl

[Rbiessy]

Are you sure this is a memory leak and not just the memory pool keeping some memory allocated? I don't expect it would need much memory on the host though.

[qnixsynapse]

Should memory pool allocate 6-12GB of system memory? It doesn't sounds normal.

[Alcpz]

IIRC We have a host memory pool only for the batched mul mats that go to onemkl. I would not expect that to allocate that much memory. Also, you shoulnd't be going through that codepath if you have oneDNN available (which you should if you have the DPCPP release compiler)

[qnixsynapse]

Yes indeed. Also, I found out that turning on SYCL graph makes it worse.

TBH, I don't know how to check for issues like this. Both Valgrind and ASan doesn't like SYCL.

[Rbiessy]

Alright thanks, we'll take a look at this test on our side too.