Repository navigation
Post your Arc results #1
Replies: 13 comments 31 replies
SYCL Flash Attention on Arc 140T (Xe-LPG): oneDNN SDPA returns wrong results for quantized KV with permuted viewsSummaryOn Intel Arc 140T (Xe-LPG+, Arrow Lake-H), the SYCL backend's oneDNN SDPA path returns numerically incorrect results when the KV cache is quantized, the layout is non-contiguous ( Separately, MTP + ngram-mod speculative decoding on SYCL leaks device memory and crashes with Hardware / Software
|
| oneDNN version | ERR |
|---|---|
| oneAPI bundled (3.11) | 1.26 |
| 3.13.3 (source build) | 1.33 |
| 3.13.3 | 1.28 |
| 3.13.3 | 1.48 |
| 3.13.3 | 1.18 |
| 3.13.3 | 1.23 |
Upgrading oneDNN from 3.11 to 3.13.3 did not fix the bug. Error stays in the 1.2–1.5 range, roughly 2500× the 0.0005 tolerance.
llama-bench Results
Scope: Ternary-Bonsai 2 27B PTQ1_0 only. These numbers apply to this specific model on this specific fork. They should not be extrapolated to other models or other builds.
| KV | ONEDNN | MKL | pp512 | tg128 |
|---|---|---|---|---|
| f16 | 0 | 0 | 167.39 ± 0.89 | 7.90 ± 0.03 |
| q8_0 | 0 | 0 | 170.46 ± 1.33 | 7.85 ± 0.03 |
| q8_0 | 1 | 1 | 174.03 ± 0.06 | 7.88 ± 0.01 |
| f16 | 1 | 1 | 212.10 ± 0.84 | 7.93 ± 0.04 |
The F16 + oneDNN row is the fastest correct configuration for this model. F16 removes the quantized-KV precondition, so the broken kernel path is not taken.
The q8_0 + oneDNN row is fast but produces wrong attention output per the test-backend-ops failure above. It should not be used for correctness-sensitive work.
Server Findings
MTP memory leak
llama-server with --spec-type draft-mtp,ngram-mod crashes at ggml-sycl.cpp:2853 with UR_RESULT_ERROR_OUT_OF_RESOURCES on the first task launch. This persists with GGML_SYCL_ENABLE_VMM=0. It is separate from the FA bug.
VMM
GGML_SYCL_ENABLE_VMM=0 is required to avoid an earlier OUT_OF_RESOURCES failure at graph reserve. This aligns with the VMM disable-by-default behavior in PR ggml-org#24673 and PR ggml-org#27689.
ngram-mod behavior
ngram-mod only fires when the model echoes a long contiguous block from the prompt. In one session:
ngram-mod: #calls(b,g,a) = 3 356 2
#gen drafts = 2
#acc drafts = 2
#gen tokens = 512
#acc tokens = 301
mean acc len = 151.50
Hit rate 2/356 calls. When it fires, it carries 300+ tokens in 2 drafts. When it does not, MTP alone carries generation at ~4–7 t/s.
ngram-mod window size
| n_max | Behavior |
|---|---|
| 256 | Catches long verbatim copies; mean accepted length up to 36 |
| 48 | Catches short targeted edits; mean accepted length up to 25, acceptance up to 89% |
Neither wins universally. A middle value (128) may capture both patterns.
Conclusion
- The oneDNN SDPA path is numerically broken on Xe-LPG for quantized KV +
permute=[0,2,1,3]+kv_view=1. Matches issue SYCL: flash attention returns a silently wrong answer for a quantised, non-contiguous KV view (Arc 140T, Xe-LPG) ggml-org/llama.cpp#27769, not fixed by oneDNN 3.13.3. - F16 KV avoids the bug and is the fastest correct configuration for this model on this fork.
- MTP + ngram-mod on SYCL has a separate memory leak causing
OUT_OF_RESOURCES. GGML_SYCL_ENABLE_VMM=0is required to avoid an earlier allocation failure.
Generated with DeepSeek <3
|
Thanks for the detailed write-up, it was really useful. I reproduced the test-backend-ops failures on my B580 too: FLASH_ATTN_EXT with a quantized KV cache in the permuted layout fails from about 1K keys up, and it only passes with both oneDNN and MKL turned off. So it's not specific to the 140T. I found the cause. The test builds its KV view with a gap between tokens (it allocates twice the heads and views half of them), and the oneDNN and MKL paths work out the token stride from the head count instead of reading it. llama.cpp's real cache has no gap, so actual models weren't affected. With Bonsai 27B, q4_0 KV and flash attention, perplexity at 4K context is the same on all three attention paths (oneDNN 8.727, MKL 8.731, the native TILE kernel 8.733), and needle-in-a-haystack recall is exact at 32K, 64K and 120K. I've added a guard so a layout with a gap falls back to the correct kernel. The full FLASH_ATTN_EXT suite now passes (4137/4137), plus a new sweep of the real cache layout from 1K to 8K keys. That fix will be in the next push. F16 KV is still a fine choice on the 140T if you have the memory for it. On the MTP crash, I suspect memory rather than a leak. The draft context needs its own KV and compute buffers, and 12 GB of shared memory gets tight fast. Could you post your exact llama-server command? On the B580 I use LLAMA_ARG_SPEC_DRAFT_UBATCH=512 plus -ctkd q4_0 -ctvd q4_0 for the draft context, and on the 140T I'd start at -c 16384 and work up from there. Also, the ternary XMX path only runs on Xe2 (B-series, Lunar Lake). On the 140T it switches itself off, so around 8 t/s is about what that memory bandwidth allows for a 27B model. |
Is it possible to run the ternary XMX for Xe-LPG+? |
|
Good question. Xe-LPG+ does have XMX (it's the main thing Arrow Lake-H added over Meteor Lake), so it's possible in principle, but not with the current code. The ternary and attention kernels are written for Xe2's XMX: 16-lane sub-groups and Xe2's 2D block loads. Xe-LPG+'s XMX comes from Alchemist (8-lane DPAS), so it would need separate kernel variants, and I don't have an Xe-LPG+ machine to build or test them on, so it's not something I'm planning. That's why the path switches itself off on the 140T. A correction to what I said before, too: 8 t/s isn't the bandwidth ceiling. If you get 10-12 t/s from a ~16 GB Q4_K_M 27B, Bonsai at 5.85 GB should do much better than 8, so the plain SYCL kernels are leaving a lot on the table on the 140T. Worth trying the Vulkan build of this branch with the same Bonsai file (steps in the guide under "Build and run (Vulkan)") and posting both numbers here. One heads-up: PrismML lists PTQ1_0 on Intel Vulkan as having been slow and occasionally hanging after about 1,900 tokens, so if it stalls, that's the likely cause. If you or anyone with Xe-LPG+ hardware wants to have a go at the XMX side, the Xe2 check is ggml_sycl_device_is_xe2() and the kernels are in ggml/src/ggml-sycl/. I'm happy to review a PR. |
|
disabling VMM made it possible to run the model whereas before it didn't even start hence the crashes, I would like to see and try to fix the SYCL kernel but i don't know where to start since i don't know, maybe if you could guide me on the start, i could do something or find something. |
|
Happy to point you in. One thing first: for generation speed, XMX probably isn't where the win is. Generating a token reads every weight once, so it's limited by memory bandwidth, and XMX mainly helps prompt reading and MTP draft checks. On the 140T the bigger opportunity is the plain PTQ1_0 mat-vec kernel it already uses, which is also much easier to work on. Where to start:
If you do want to try XMX later, the gate is ggml_sycl_device_is_xe2() in ggml/src/ggml-sycl/ptq1-t2.cpp (line 67), and those kernels live in ptq1-t2.cpp and ternsycl/. They'd need 8-lane DPAS variants for Xe-LPG+. I can't test on that hardware, but I'm happy to review whatever you come up with. |
|
System info: Context size:
In most local models I have tried, this task end up getting brute forced and often is in an infinite loop. Result from the Terminal:
|
|
Thanks for testing, and for including the full log. 55 t/s sustained across almost 6,000 generated tokens on Windows is a solid result.
The "4 tokens" isn't wrong, it's prompt caching. Your prompt was 381 tokens and 377 of them were already in the server's cache from the previous turn (f_sim_best = 1.000), so only 4 new ones needed processing. That conversation used about 6,350 tokens of context (n_tokens = 6353). The limit is set by -c, which is 131,072 (128K) in run-bonsai.bat, and the server prints it at startup as n_ctx.
A couple of questions so I can read the numbers properly: did you use run-bonsai.bat as is, or your own command? The .bat starts with thinking turned off, so the model does its working in the answer itself. And did it get the timer puzzle right (it's 9 minutes: start both, flip the 7 when it runs out at 7, then flip it again when the 4 runs out at 8, and it finishes at 9)?
|
|
Thanks for the explanation. The puzzle is particularly hard for the AI models. Gemma, Ornith, Ling, Laguna, those all failed the task, but were great at many other tests I tried with similar difficulty. Ternary-Bonsai end up the answering with this:
While it didn't produce the correct answer, it did think of the correct combo for the solution (4+4+1) Flip both, 4 runs out flip again, 7 runs out - flip, 4 runs out - flip the 7 again as 1 min is on the other end.
With the flag
|
|
Thanks for running both, that's a really useful comparison. Thinking on is the right setting for puzzles like this. It's slower per token (42 vs 55 t/s, since fewer of the drafted tokens get accepted while it's reasoning), but it needed about 1,000 tokens instead of 6,000, so the answer actually came back about 4x sooner. The .bat starts with thinking off because that's faster for coding, but I'll add a note to the guide about switching it on for reasoning tasks. Nice to hear it beat Gemma and the others on this one.
|
|
Tested the Windows System:
Builds:
I tested a regular dense GGUF rather than Bonsai first: Gemma 4 31B QAT Q4 Settings:
Results on the same B60:
Decode was essentially identical:
So on the B60, long-context decode performance for this regular Q4 dense model is effectively the same as my newer SYCL build, while prompt processing is dramatically faster. I also tested Gemma 4 MTP speculative decoding also works. I tested the QAT Q4 model with:
Reasoning budget works as well. For Gemma 4 I needed I haven't run Bonsai 2 on the B60 yet, but that's next. |
|
Thanks for testing on the B60, and with Gemma 4 31B too. Good to see the prompt-side kernels carry over to a non-Bonsai model, and that its MTP drafting works. |
|
arc-b580 on 258v (Arc 140v): crashes on every model
|
|
Thanks for the report, that's a real bug on my side. Two of the XMX kernels use an IGC builtin ( |
|
The fix is in: f31a4f6 on |
|
Quick update: v0.6 still had other builtins of that kind in the Bonsai matrix path (the ternary DPAS and 2D block loads), so it may still fail the same way on the 140V. v0.7 removes all of them: https://github.com/Torchit1/llama.cpp/releases/tag/arc-b580-v0.7. If you get a chance, that's the one to retest. |
|
Thanks for the update. It works more or less the same as the local fix that made by Gemini. I got about 210 t/s pp and 12.7 t/s tg (without MTP) on 140v |
|
Hi. Created this github account to say thank you for this fork. I've tested on B570 (linux 7.2.9, gentoo). (Each test was performed at least 2 times in a row.) For Gemma4, llama-bench is 755/36 (p/g) for kv q8_0, 755/38 for kv q4_0, 875/37 for f16, against 230/38 for kv f16 SYCL on llama-cpp 0.6.0 (pp improved significantly). Token generation stays at 55-65 with MTP when using web-ui. As for Tern-Bonsai, it will not run on 0.6.0, so I can't compare directly - but both prompt proccessing and token generation is better than "standard" Qwen-3.8-27B on any quantization: |
|
Adding Bonsai2-27B multi-session results on a 24 GB Arc Pro B60. Hardware/software:
These are server measurements with deterministic synthetic repository inputs. The allocated window was 3 × 262,144 input-plus-output tokens. Each request actually contained 258,048 input tokens and generated 1,024 tokens, with all three requests active concurrently.
The baseline used q4_0 KV for target and draft, MTP4, batch 2048/ubatch 1024, draft ubatch 512, three non-unified slots, compact recurrent replay and no RAM archive. The compact main-model recurrent allocation was 570.586 MiB versus 948.375 MiB padded: 377.8 MiB saved. The patch retains the arithmetic kernels; an isolated speed gain from compaction was not established. I published the patches, exact environment/launch settings, pinned build recipe, benchmark generator and selected raw measurement records: https://github.com/jc1122/bonsai-b60-repro. After its build steps, A separate eviction/resume experiment on 18 October used The full-context test demonstrates fit/concurrent generation; coding accuracy at 256K remains unmeasured. The VRAM margin is tight and depends on this pinned stack. Credit for the ternary/XMX, attention and speculative-decoding foundation belongs to Torchit, Intel and the upstream Bonsai work. |
|
Thanks for this, it's a really thorough write-up. Three 256K sessions on a 24 GB card is further than I'd pushed it, and the resume result (158 s down to 0.7 s) is great to see.
The compact recurrent allocation looks useful for the 12 GB B580 too. Would you be happy for me to fold that idea into the next release, with credit to you?
|
|
can you check the fork i did, since you asked on the reddit, pr 294 is working sycl on Arc 140T baseline is 50 t/s, improved is 100-110 t/s, and experimental branch which went up to 170 but outputs nonsense |
|
Thanks, I'll take a look when I can. What speeds are you getting? |
|
Had a look at your fork. The pr-294 tuning for the 140T looks really promising: 50 → 100–110 is a big jump. Worth opening that as a PR on PrismML, they haven't got many Arrow Lake numbers. On xmx-experiment, I think I can see the garbage. Does it pass test-backend-ops with that change? |
|
I'm lazy XD Driver:Intel® Arc™ 140V GPU (16GB) 32.0.101.9150 You can find card, driver, OS👆.Plus what computer I'm using. context size is inside the bat(131072) I runned the bat,and got: |
|
Feel free to give me some new tests(Though I might be too lazy to be quick XD) |
|
Thanks, this is great, and the first Lunar Lake + Windows result I've seen for this. 10.6 t/s on the 140V lines up with the memory bandwidth, so it looks healthy.
If you get a minute, could you run the same bat once more with set GGML_SYCL_PTQ1_T2=all at the top? That turns on the TernSYCL kernels, and on my B580 they're 20-30% faster for generation. Curious if Lunar Lake sees the same.
|
|
Only feels like it's even slower...... BTW Lenovo Computer Manager provides three modes（Translated by doubao）: |
|
Thanks for running it, that's really interesting. It's the opposite of the B580, where T2 is 20-30% faster. Could be the iGPU sharing memory bandwidth, or the power mode like you said. If you get a chance, could you try it once on the Lenovo "Performance" mode? That'd tell us if it's thermals or the kernel itself. Either way this is useful for the PrismML folks, they haven't had Lunar Lake numbers. |
|
OK * 2 |
|
This time,I just tested six combinations. Result-NoT2-balanced.txt |
|
btw I use 19.9GB(63%) sharing memory. |
|
Thanks for running all six, this is really useful. So it was the power mode, not the kernel: in Performance, T2 gives +7% generation (11.1 → 11.9) and prompt goes 274 → 359. On Balanced/Eco the laptop's power limit makes T2 noisy.
On the 19.9 GB: the model itself is about 5.9 GB, and most of the rest is the 131K context. If you don't need that much, -c 32768 will cut the memory a lot and leave more bandwidth headroom on shared memory.
I'll add your 140V numbers to the results table at the top, credited to you.
|
|
So sorry @Torchit1 |
I run the bench randomly just now,and surprisingly found that NoT2 in performance mode run faster than before on pp512,and slower on tg128??? Result-NoT2-performance-another1.txt |
|
I can't exactly remember whether I zipped the model.Anyway,There's no improvement on T2 after unzip. |
|
No worries at all, thanks for re-running everything! With the uncompressed runs the picture is clearer: on the 140V, T2 speeds up prompt reading by about 1.3-1.5x (roughly 235-260 → 340-360 pp512), while generation is about the same (~11.5-11.8 without vs ~11.7-11.9 with). That makes sense: T2 repacks the weights into a 2-bit layout (about 30% more bytes to read per token), and on a shared-memory iGPU generation is limited by bandwidth. So I was wrong earlier about +7% generation, sorry. For long prompts T2 on is the clear win; for short chats it's roughly a wash. And good call on NTFS compression: keep the model folder uncompressed. Windows has to decompress memory-mapped files on the fly, which would explain the run-to-run swings. |
|
OK! |
|
Dual B580s here. Just recently got my second B580, I had an AMD 6700 XT next to the other B580 before. These speeds are like a generational leap for me. And it's like 150W less in power. I couldn't get the Ternary Bonsai model to work that's included in the run-bonsai.bat in the v0.10.1 release for some reason, it complains about being out of resources: As for my daily driver model (Unsloths Qwen 3.8 27B Q4_K_S) the results are amazing. I found that on longer prompts to process it stretches its legs a bit more, giving more representative numbers than pp512. I'm still going back and forth between -ub 1024 and 2048, 1024 gives more flexibility for longer context windows with imo not that big a trade off in prompt processing speeds. The tests are done on Windows 11 25H2, AMD Ryzen 7 3700x. One of the B580s (Intel Limited Edition) is in the main PCIe 4.0 X16 slot (card is wired to X8), the other B580 (Asrock Challanger) is in the second PCIe 4.0 X16 slot, wired to X4 (don't think it makes too much of a difference here). Comparing this to the Vulkan version of the latest llama.cpp build it's easy to see Vulkan starting with a higher prompt processing speed, but dropping of way quicker than SYCL, giving less than half the speed at 35k tokens. |
|
Thanks, this is a great comparison, and the long-prompt result is exactly what I was hoping to see: SYCL holding up at 16K+ while Vulkan falls off. Generation being a touch behind Vulkan is known; I'm working on that right now (the GPU sits idle between steps, and I'm moving decode onto SYCL command graphs to fix it). The Bonsai error is almost certainly a multi-GPU bug in the ternary path that a dual-B70 tester hit too. It's fixed on the |
|
Yep, but we've already seen how that performs. Completely off topic, but it might be useful for people with multiple Intel GPUs on Windows: I'm not sure if it's a driver feature or a Windows behavior, but about a minute after the GPUs become idle the secondary GPUs evacuate all their VRAM contents into system RAM, then load it back in when work resumes - apparently this happens with AMD GPUs too. This is especially annoying when a harness is waiting for a compile to finish and you have limited system RAM already. I've been looking at this PR, and saw this comment which keeps pinging a SYCL device every 10 seconds so it doesn't go to "sleep". I adapted it for my needs (enumerates all SYCL devices, not just the first one), and can say it works well as an "accessory" for the inference engine. #include <sycl/sycl.hpp>
#include <chrono>
#include <thread>
#include <iostream>
#include <vector>
int main() {
std::vector<sycl::device> gpus;
for (const auto& d : sycl::device::get_devices()) {
if (d.has(sycl::aspect::gpu) || d.has(sycl::aspect::accelerator)) {
gpus.push_back(d);
}
}
if (gpus.empty()) {
std::cerr << "No SYCL GPU/accelerator devices found.\n";
return 1;
}
std::vector<sycl::queue> queues;
for (const auto& d : gpus) {
std::cout << "Keeping alive: "
<< d.get_info<sycl::info::device::name>()
<< " [" << d.get_info<sycl::info::device::vendor>() << "]"
<< std::endl;
queues.emplace_back(d);
}
std::cout << std::flush;
while (true) {
for (auto& q : queues) {
q.submit([&](sycl::handler& h) {
h.single_task([] {
// no-op kernel
});
});
}
for (auto& q : queues) {
q.wait();
}
std::this_thread::sleep_for(std::chrono::seconds(10));
}
}Can be compiled with the oneAPI base toolkit. |
Uh oh!
There was an error while loading. Please reload this page.
Uh oh!
There was an error while loading. Please reload this page.
Running the arc-b580 branch? Post your card, driver, OS, context size and the numbers you get (llama-bench tg128/pp512, or the server's t/s). Results from other setups help most.
Community results
Thanks to everyone who's tested this. My own B580 numbers are in the README. Bonsai 2 27B (PTQ1_0), llama-bench pp512 / tg128 unless noted.
Got a card that's not here (A-series, B50, B65, or more B70s)? Post your numbers below and I'll add them.
All reactions