sycl: single-device still requires sync

[malsbat]

Overview

Fix garbage output generated on second prompt in multi-turn conversation with model Qwen3-14B-Q4_K_M.

Additional information

This can be reproduced with:

$ ./build/bin/llama-cli -m /models/Qwen3-14B-Q4_K_M.gguf -ngl 100
> foo
# good output
> bar
# garbage output

The problematic commit, 32b741c, was identified with help from git bisect.

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: YES, help to identify issue once the right commit was found.

[malsbat]

@hmscider, please take a look at this, the repro steps are in the description. If this is the right fix let me know and I will mark it ready for review or close it if not.

[andreyzagoruy]

Can confirm garbled output for Gemma 4 12B as well

[NeoZhangJianyu]

It's good job!

Thank you!

[hibiday]

I arrived at similar code from a different (and probably incorrect) hypothesis at #25825. I confirmed that this code alone fixes the problem in Qwen3.5-0.8B.

0.02.364.580 I
0.02.364.622 I system_info: n_threads = 16 (n_threads_batch = 16) / 32 | CPU : LLAMAFILE = 1 | OPENMP = 1 | REPACK = 1 |
0.02.364.626 I perplexity: tokenizing the input ..
0.02.375.187 I perplexity: tokenization took 10.557 ms
0.02.375.229 I perplexity: calculating perplexity over 2 chunks, n_ctx=4096, batch_size=2048, n_seq=1
0.03.868.580 I perplexity: 1.49 seconds per pass - ETA 0.03 minutes
[1]7.5567,[2]7.1586,
0.04.596.944 I Final estimate: PPL = 7.1586 +/- 0.30513

| model                          |       size |     params | backend    | ngl |  fa |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --: | --------------: | -------------------: |
| qwen35 0.8B Q8_0               | 763.78 MiB |   752.39 M | SYCL       |  99 |   1 |          pp2048 |     11506.75 ± 25.84 |

build: bbe884c-pr25741 (9925)

[meatposes]

We hit this same bug in production (Arc Pro B70, Qwen3.6-27B Q4_K, 131k ctx) and root-caused it
while this PR was open. The sync does fix the symptom, but I believe the actual defect is
narrower, and fixing it directly lets you keep the single-device fast path (and its ~6% PP)
that the original commit was trying to preserve.

Root cause: a use-after-return, not missing device synchronization.

const sycl::half scale_h = (sycl::half) (1.0f / kq_scale);   // stack local
ggml_sycl_pool_alloc<sycl::half> scbuf(ctx.pool(), 1);
stream->memcpy(scbuf.get(), &scale_h, sizeof(sycl::half));   // async, sourced from the stack

The queue is in-order and the dnnl stream wraps the same queue, so every device-side hazard
(staging kernels -> SDPA -> permute -> pool-slot reuse) already serializes correctly. The one
thing that escapes ordering is host-side: this memcpy's source is a stack local. The copy is
queued behind the K/V staging kernels; once n_kv is large enough (>= ~26k on B70 -- staging is
tens of MB) the function returns before the copy runs, the frame is recycled, and the SDPA
receives a garbage scale. Output collapses to a single repeated token and the poisoned KV cache
corrupts the rest of the session. Short contexts win the race by accident, which is why it looks
like a "second prompt" / long-context bug, and test-backend-ops never sees it (FLASH_ATTN_EXT
caps at kv=1024, and single-op runs sync between ops).

wait_and_throw() fixes this only as a side effect: blocking the host keeps the frame alive
until the copy completes. That works, but it re-pays a host round-trip on every prefill FA call,
and it leaves the UAF in place -- if anyone later re-loosens the sync to reclaim the perf, the
garbage returns.

Direct fix: upload the scale synchronously, once, into a persistent device scalar cached per
(device, value) -- the scale is constant per model. With that, the single-device no-sync path is
safe. Verified on B70: temp-0 output byte-identical to the native FA path through 32k-deep
prefill, prefill depth-flat at 820-840 t/s (vs ~350 native at 32k depth), i.e. the fix keeps the
full speed this feature was added for.

Deterministic reproducer (unmodified build, GGML_SYCL_F16=ON + GGML_SYCL_DNN=ON): send a single
~28k-token prompt ending "Reply with exactly: HELLO WORLD" at temperature 0.
GGML_SYCL_FA_ONEDNN=1 -> repeated-token garbage / server 500; =0 -> "HELLO WORLD". With the
scale fix applied and the sync still removed, both answer correctly.

Happy to open the alternative as a follow-up PR (patch is +46/-10: the scale fix, the multi-GPU
wait kept unchanged, an optional GGML_SYCL_FA_ONEDNN_MAX_KV ceiling, and large-KV
test-backend-ops cases covering the kv=1024 blind spot) -- or fold any part of it into this one
if @malsbat prefers.

[meatposes]

Follow-up PR with the root-cause fix is up: #25880

[malsbat]

Thanks @meatposes for the detailed write-up. I only wanted to highlight the issue and suggest a possible fix in this PR. I'm happy to abandon this for your solution which doesn't lose the performance benefit. I'll defer to @arthw for how he'd prefer to handle this.

[arthw]

@malsbat
If #25880 can fix this issue without adding more wait(), I prefer to use it.
Then this PR can be closed.

Thank you very much! :)

[malsbat]

Closing, #25880 fixes the issue for me also.