Repository navigation
sycl: remove unused flash attention kv buffers - #27689
Conversation
github-actions
Bot
added
ggml
|
Hi @Pratyush-gg, thanks for your contribution! Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:
Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below. |
arthw
left a comment
arthw
left a comment
There was a problem hiding this comment.
It's good job!
Summary main comments:
-
remove the warning of get_memory_info()
I will create a PR to refactor this part and remove the warning in most cases.
So, no need to update it in this PR. -
Usage of GGML_SYCL_DISABLE_DNN / GGML_SYCL_DISABLE_VMM .
There are env vars for DNN and VMM: GGML_SYCL_ENABLE_DNN / GGML_SYCL_ENABLE_VMM.
Do you want to define new vars named as GGML_SYCL_DISABLE_DNN / GGML_SYCL_DISABLE_VMM?
If yes, please use another names; update the SYCL.md to explain them like GGML_SYCL_ENABLE_DNN; need to print the value after read them.
Thank you!
| #if (defined(__SYCL_COMPILER_VERSION) && __SYCL_COMPILER_VERSION >= 20221105) | ||
| if (!has(sycl::aspect::ext_intel_free_memory)) { | ||
| std::cerr << warning_info << std::endl; | ||
| if (!warned) { |
There was a problem hiding this comment.
This changed code want to print out the warning once.
But because this variable is static, it only prints out once even if there are multiple GPUs.
Get the free memory of GPU is always an issue.
I want to update it by a new PR to fix it and remove the warning.
I suggest not changing this part in the PR.
| g_ggml_sycl_enable_optimize = ggml_sycl_get_env("GGML_SYCL_ENABLE_OPT", 1); | ||
| g_ggml_sycl_enable_graph = ggml_sycl_get_env("GGML_SYCL_ENABLE_GRAPH", 0); | ||
| g_ggml_sycl_enable_dnn = ggml_sycl_get_env("GGML_SYCL_ENABLE_DNN", 1); | ||
| g_ggml_sycl_enable_dnn = ggml_sycl_get_env("GGML_SYCL_ENABLE_DNN", 1) && !ggml_sycl_get_env("GGML_SYCL_DISABLE_DNN", 0); |
There was a problem hiding this comment.
| g_ggml_sycl_enable_dnn = ggml_sycl_get_env("GGML_SYCL_ENABLE_DNN", 1) && !ggml_sycl_get_env("GGML_SYCL_DISABLE_DNN", 0); | |
| g_ggml_sycl_enable_dnn = ggml_sycl_get_env("GGML_SYCL_ENABLE_DNN", 1) ; |
GGML_SYCL_DISABLE_DNN is renamed to GGML_SYCL_ENABLE_DNN in base code.
| g_ggml_sycl_fa_onednn = ggml_sycl_get_env("GGML_SYCL_FA_ONEDNN", 1); | ||
| g_ggml_sycl_fa_onednn_max_kv = ggml_sycl_get_env("GGML_SYCL_FA_ONEDNN_MAX_KV", 0); | ||
| g_ggml_sycl_enable_vmm = ggml_sycl_get_env("GGML_SYCL_ENABLE_VMM", 1); | ||
| g_ggml_sycl_enable_vmm = ggml_sycl_get_env("GGML_SYCL_ENABLE_VMM", 1) && !ggml_sycl_get_env("GGML_SYCL_DISABLE_VMM", 0); |
There was a problem hiding this comment.
| g_ggml_sycl_enable_vmm = ggml_sycl_get_env("GGML_SYCL_ENABLE_VMM", 1) && !ggml_sycl_get_env("GGML_SYCL_DISABLE_VMM", 0); | |
| g_ggml_sycl_enable_vmm = ggml_sycl_get_env("GGML_SYCL_ENABLE_VMM", 1); |
GGML_SYCL_DISABLE_VMM is renamed to GGML_SYCL_ENABLE_VMM in base.
| case GGML_TYPE_Q1_0: | ||
| case GGML_TYPE_Q4_0: | ||
| case GGML_TYPE_Q8_0: | ||
| case GGML_TYPE_Q2_K: |
| static bool ggml_sycl_supports_reorder_esimd(enum ggml_type type) { | ||
| #ifdef GGML_SYCL_DMMV_HAS_ESIMD | ||
| switch (type) { | ||
| case GGML_TYPE_Q2_K: |
|
I will verify the function on iGPU. |
|
@Pratyush-gg Log: |
|
@arthw I have made the changes as suggested by you. I have updated get_scratchpad_mem() to return dnnl::memory(scratchpad_md, eng, nullptr) when scratchpad_md.get_size() == 0. I have also removed the DISABLE_* variable aliases in favor of the standard GGML_SYCL_ENABLE_* flags. I was using these setting when running for igpu before: Could you please retest on discrete and igpu with the latest commit? Thanks!" |
|
@Pratyush-gg Don't use it I'm afraid the issue is not easy to fix. Thank you! |
|
@arthw thanks for testing this again this OOM in I tried to reproduce it on my own hardware my i9-14900K's UHD 770 iGPU, and separately a laptop with Iris Xe using the same model, flags, and env vars you posted ( The one structural difference I can see between our setups is device count. I only have a single Level-Zero GPU on either machine, while your Could you help narrow it down with two more tests when you get a chance?
If you can reproduce it again, a trace would help pin down the actual error under the generic OOM message: |
|
I had a similar OOM with a Intel Meteor Lake iGPU and B580 eGPU. Attempting to run on the MTL iGPU resulted in this. Hiding the B580 with export ONEAPI_DEVICE_SELECTOR=level_zero:1caused the model to load normally. Ran a trace, file is attached. Options were |
|
@Pratyush-gg I'm checking this issue: crash on iGPU only. How do you think? Thank you! |
|
@arthw agreed lets hold this until its sorted. Worth noting your dgpu worked fine before this PR and broke with an earlier version of it, so I won't assume the current crash is unrelated to my changes either. Better to find the real cause first. @KellenRenshaw Could you send a smaller trace? The full ze_debug.txt cut off before the actual crash. If you run: that gives you the line number of the crash, then pull a few thousand lines before it with sed and attach that instead. Should land right on the failing call. Thanks both I'll keep testing on my end too. |
|
@Pratyush-gg Sorry about that, didn't realize it did that. Attaching a full log. |
|
@KellenRenshaw looks like GGML_SYCL_DEBUG was on but SYCL_UR_TRACE and ZE_DEBUG werent. Maybe llama server runs as a systemd service here, exporting those in your shell dont work, systemd services get their own environment. Could you try and add them to the test instead? Add: Then: |
|
Ok, pulled it out of the systemd unit, some sort of buffering was interfering with the output. Here's the run configuration: source /opt/intel/oneapi/setvars.sh
:: initializing oneAPI environment ...
-bash: BASH_VERSION = 5.3.9(1)-release
args: Using "$@" for setvars.sh arguments:
:: ccl -- latest
:: compiler -- latest
:: debugger -- latest
:: dev-utilities -- latest
:: dnnl -- latest
:: dpl -- latest
:: mkl -- latest
:: mpi -- latest
:: pti -- latest
:: tbb -- latest
:: tcm -- latest
:: umf -- latest
:: oneAPI environment initialized ::
export ZES_ENABLE_SYSMAN=1
export SYCL_UR_TRACE=2
export ZE_DEBUG=1
export GGML_SYCL_USM_SYSTEM=1
export LD_LIBRARY_PATH=/opt/llama.cpp/SYCL-mtl/lib:$LD_LIBRARY_PATH
/opt/llama.cpp/SYCL-mtl/bin/llama-server --models-preset /home/kellen/models/models-mtl.preset --port 8081 | tee ze_debug.txt |
|
Quick note, setting GGML_SYCL_USM_SYSTEM=0 didn't change the behavior. |
| } | ||
| dnnl::memory get_scratchpad_mem(const dnnl::memory::desc & scratchpad_md, | ||
| const dnnl::engine & eng, const queue_ptr q) { | ||
| if (scratchpad_md.get_size() == 0) { |
There was a problem hiding this comment.
- My PR sycl : fix oneDNN scratchpad breaking the pool free order #28704 removes this function entirely. Whoever merges second should remove the function on conflict resolution.
- My change also passes nullptr in case of 0 size scratchpad, so nothing you fixed here is lost.
- The LIFO assert which @arthw reported hits in a second scenario when
get_size() > 0and VMM works. The scenario is described in Eval bug: SYCL crash in ggml_sycl_pool_vmm::free - oneDNN scratchpad breaks LIFO pool order #28660 and I verified that the change here does not fix it (all fine, you had a different case). - You may want to change
std::max((size_t)1, ...)in description as your patch uses an early return now. - I have machine with Arc Prov B70 and UHD Graphics 770 visible through LevelZero. Do you need help?
|
@Pratyush-gg , the OOM errors which were met here are not related to your changes, but are caused by #28765. I managed to reproduce them and found the root cause on a clean master branch. Your code is clean as far as this crash is concerned. Your hypothesis that this might be something that only surfaces with 2+ level zero devices alive in the process was very close to the truth, but the condition for the issue is having a non-zero SYCL index for the device. @arthw, hypothesis about crash on iGPU only is false. The issue hits whatever GPU has non-zero index. You can verify that by swapping your devices with Workaround for the issue is to set I have a working fix for #28765 based on using the correct device index instead of the hardcoded 0 in |
|
@lslusarczyk Thank you! |
@arthw , @KellenRenshaw , @Pratyush-gg , the PR with OOM fix is ready here #28895. You are welcome to verify and review. |
|
@Pratyush-gg
Could you rebase this PR? Thank you! |
0a1554f to
7d7627d
Compare
Pratyush-gg
changed the title
lslusarczyk
left a comment
lslusarczyk
left a comment
There was a problem hiding this comment.
Dead code removal is done perfectly. @arthw, go ahead with merging this
arthw
left a comment
arthw
left a comment
There was a problem hiding this comment.
It's good job!
I have tested with UT and benchmark.
All are passed.
Thank you!
arthw
added
the
merge ready
(cherry picked from commit 5e5b628)
5 participants
Overview
Remove the unused ggml_sycl_fattn_kv_buffers implementation and related code. K/V F16 conversion now uses extra K_buffer_ptr / extra.V_buffer_ptr directly with assertions. These buffers became unreachable after #27629, so the fallback allocations were dead code.
Additional information
Iris Xe, Windows 11, oneAPI 2026.1
FLASH_ATTN_EXT: 350/350 passed for q4_0; 363/364 for q8_0 (known #27769 failure)
llama-bench: Qwen2.5-1.5B Q4_K_M, -fa 1, q8_0 KV cache ran successfully
Requirements
(the warning i am getting is maybe because I used AI to format the text properly)