SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt proc… by johnkarlhill · Pull Request #25025 · ggml-org/llama.cpp

johnkarlhill · 2026-06-26T01:24:12Z

Adds a flash attention path that routes Q·K^T and S·V matrix multiplies
through oneMKL GEMM, enabling XMX hardware acceleration on Intel GPUs.

Motivation

The existing SYCL flash attention kernels (VEC, TILE) run entirely in
SYCL subgroup operations. On Intel Arc GPUs with XMX matrix engines
(Battlemage and later), oneMKL GEMM can process the large matmuls in
attention significantly faster — particularly at high context lengths.
~~where the KV cache is quantized.~~

When it activates

~~- KV cache is quantized (q8_0, q4_0, q4_1, q5_0, q5_1, or any K-quant)~~

Fast Attention is enabled (--flash-attn on or -fa)
K sequence length ≥ 1024 tokens (covers the full --batch-size)
~~- Q sequence length ≥ 128~~
Q sequence length ≥ 32 (routes all multi-token prefill through MKL
while keeping single-token decode and MTP spec drafts on the existing
TG-optimized VEC kernel)

~~These thresholds route prompt processing through MKL while leaving
single-token decode to the existing TG-optimized kernels. ~~
~~The path is never activated for f16/bf16 KV cache — those already perform well with the TILE kernel and graph capture.~~

All KV cache types benefit from XMX acceleration — F16 (the default),
BF16, F32, and quantized (q8_0, q4_0, q4_1, q5_0, q5_1, K-quant).
The MKL kernel converts non-F16 K/V to F16 via to_fp16_sycl before
GEMM, so no additional conversion code was needed.

Set GGML_SYCL_ENABLE_MKL_FA=0 to force the TILE/VEC path for A/B
testing or power comparison.

Implementation

All logic is in one new file, fattn-mkl.cpp (567 lines). The pipeline:

Dequantize K/V to fp16
For each KV head: pack all GQA query heads into a single fp16 buffer
Chunked KV loop (8192-token chunks):
- MKL GEMM: KQ = Q_batched × K_chunk^T
- Online softmax SYCL kernel (row-wise, with running max/sum)
- MKL GEMM: VKQ_chunk = S × V_chunk
- Accumulate: VKQ_accum += VKQ_chunk
Normalize each GQA head by KQ_sum and scatter to output

GQA groups sharing a KV head are batched into single GEMM calls —
6 query heads × 1020 tokens = 6120 rows in one MKL call, amortizing
launch overhead.

All benchmarks: Qwen 3.6 27B UD Q5_K_XL, MTP enabled.

Context	KV Cache	PP t/s	TG t/s
32K	f16 (default)	~671	~15
32K	bf16	~668	~12
32K	q8_0	~671	~15
110K	q8_0	~335	~17

~~For comparison, stock bf16 KV cache + FA off on the same GPU achieves 822 t/s PP at 8K — the MKL path with q8_0 is within 1% while using quantized memory.~~
The f16/bf16 numbers are with FA-on and the MKL path — matching the
quantized-cache baseline that previously required --cache-type-k q8_0

Testing

test-backend-ops: 3605/3605 FLASH_ATTN_EXT tests pass (all
quant types, head sizes 64–512, causal/non-causal masks, sinks,
max_bias, GQA ratios, multi-batch)
Multi-batch: parallel-2 at 32K context, stable throughput,
no coherence errors
Cache reuse: full-context slot restore with LCP similarity,
MKL path correctly handles the reprocess delta
Build isolation: all code is SYCL-only, gated behind
BEST_FATTN_KERNEL_MKL enum; other backends and non-quantized
paths are completely unaffected

Known limitations

No graph capture: MKL GEMM's internal queue management is
incompatible with SYCL command graph replay. The existing
GGML_SYCL_DISABLE_GRAPH default (1) handles this.
No ALiBi: max_bias == 0.0f asserted; models needing ALiBi
will fall through to the TILE kernel.
No sinks tensor: dst->src[4] is not yet supported.

Debug output

~~Timing instrumentation is gated behind MKL_FA_DEBUG=1. In normal operation the MKL path produces no output.~~
Timing instrumentation is gated behind GGML_SYCL_MKL_FA_DEBUG=1.
Output fingerprint for correctness verification is available with
GGML_SYCL_MKL_FA_DIAG=1. In normal operation the MKL path produces
no output.

AI disclosure

Claude Code was used for SYCL boilerplate (ND-range kernel launches,
ggml_sycl_pool_alloc patterns) and initial drafting of the chunked KV
loop. All algorithmic decisions — oneMKL GEMM integration, online
softmax with GQA batching, activation thresholds, chunk sizing — were
human-directed. Comprehensive testing (3605 test-backend-ops,
multi-quant and multi-batch coherence validation, performance
benchmarking at contexts up to 110K) was performed manually.

🤖 Generated with Claude Code using DeepSeek-V4-Pro

ggml-gh-bot · 2026-06-26T01:28:39Z

Hi @johnkarlhill, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

AI-generated content: This project does not accept PRs, descriptions or commit messages that are fully or predominantly AI-generated. If you have used AI to assist you in writing code, please make sure to disclose that explicitly.

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

arthw

@johnkarlhill

It's good to see this PR to enable XMX in FA.

Could you share which LLM show good performance increasing by this PR?
I use Qwen3.6 and can't trigger the oneMKL path on FA.

Thank you!

johnkarlhill · 2026-06-26T03:19:22Z

Adding before and after... both compiled with arch flags to show side-by-side. Compiling without arch flags will degrade performance from these numbers but should still be better than stock.
PR25025 - Qwen3.6-27B-MTP-UD-Q5_K_XL on B70.txt
b9752 - Qwen3.6-27B-MTP-UD-Q5_K_XL on B70.txt

I'll add more models if needed.

arthw · 2026-06-26T07:36:25Z

@johnkarlhill
1.
Could you provide a smaller LLM case to show the perf increase for this PR?
Including the whole cmd.

For user, how to trigger the new code in usage?

Thank you!

maxious · 2026-06-26T09:05:24Z

Tested on dual Intel Arc Pro B60 (Battlemage, 24GB each), oneAPI 2026.0, MKL 2026.0, targeting bmg-g31 AOT.

Short context (pp512 — MKL path NOT active)

No regression vs master:

Model	Size	Master pp512	PR pp512
gpt-oss 20B Q8_0	11.3G	854	851

Long context (pp2048 — MKL path active)

Model	Size	Master pp2048	PR pp2048	Delta
llama-2 7B Q2_K	2.6G	950	1,102	+16%
Llama-3 8B Q4_0	4.3G	888	968	+9%
gpt-oss 20B Q8_0	11.3G	503	504	0%
Qwen3.6-35B-A3B MoE Q3_K	12.8G	575	575	0%

Decode speeds (tg128) unchanged across all models (±2%).

Full commands

# Master
./build-master/bin/llama-bench -m model.gguf -p 512,2048 -n 128 -ngl -1 -fa 1

# PR #25025
./build-pr/bin/llama-bench -m model.gguf -p 512,2048 -n 128 -ngl -1 -fa 1

johnkarlhill · 2026-06-26T11:45:15Z

Updated arch table:

GPU	-DGGML_SYCL_DEVICE_ARCH
Arc A770 / A750	acm_g10
Arc A580	acm_g12
Arc A380 / A310	acm_g11
Arc B580 / B570 / Pro B70	bmg_g21
Flex / Data Center Max	pvc
Integrated (Meteor Lake)	mtl_u
Integrated (Lunar Lake)	lnl_m

All of these use underscores (_), never hyphens (-).

I incorrectly listed Arc A770 / A750 as acm_g12.

"For user, how to trigger the new code in usage?"
Use "--batch-size N" where N is a value >= 1024.

jlionhan · 2026-06-26T16:35:39Z

I hope this helps.

255H, Arc 140T, 32GB RAM, llama-cli

Build options:

cmake --fresh -B build -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DCMAKE_BUILD_TYPE=Release -DGGML_SYCL=1 -DBUILD_SHARED_LIBS=0 -DGGML_SYCL_F16=1

after PR:

model	size	params	backend	ngl	threads	type_k	type_v	fa	test	t/s
gemma4 26B.A4B Q5_K - Medium	17.80 GiB	25.23 B	SYCL	99	6	q8_0	q8_0	1	pp512	176.84 ± 3.13
gemma4 26B.A4B Q5_K - Medium	17.80 GiB	25.23 B	SYCL	99	6	q8_0	q8_0	1	pp1024	202.22 ± 2.48
gemma4 26B.A4B Q5_K - Medium	17.80 GiB	25.23 B	SYCL	99	6	q8_0	q8_0	1	pp2048	212.19 ± 1.95
gemma4 26B.A4B Q5_K - Medium	17.80 GiB	25.23 B	SYCL	99	6	q8_0	q8_0	1	tg128	12.13 ± 0.22

before PR:

model	size	params	backend	ngl	threads	type_k	type_v	fa	test	t/s
gemma4 26B.A4B Q5_K - Medium	17.80 GiB	25.23 B	SYCL	99	6	q8_0	q8_0	1	pp512	178.96 ± 6.26
gemma4 26B.A4B Q5_K - Medium	17.80 GiB	25.23 B	SYCL	99	6	q8_0	q8_0	1	pp1024	165.24 ± 3.03
gemma4 26B.A4B Q5_K - Medium	17.80 GiB	25.23 B	SYCL	99	6	q8_0	q8_0	1	pp2048	139.39 ± 1.68
gemma4 26B.A4B Q5_K - Medium	17.80 GiB	25.23 B	SYCL	99	6	q8_0	q8_0	1	tg128	12.02 ± 0.38

build: f728ada (9793)

Thank you.

However, I am observing intermittent behavior where the most recently entered prompt is not being processed, and the model instead generates a response to the previous prompt. I believe further testing is needed to confirm whether this issue is reproducible and to identify the underlying cause. This behavior may be unrelated to this PR.

johnkarlhill · 2026-06-26T22:26:46Z

However, I am observing intermittent behavior where the most recently entered prompt is not being processed, and the model instead generates a response to the previous prompt. I believe further testing is needed to confirm whether this issue is reproducible and to identify the underlying cause. This behavior may be unrelated to this PR.

I can reproduce the behavior on Gemma4 models and working on a fix. This behavior does not exist on Qwen models. Multiple folks have tested with a few different Qwen models and can't reproduce this. It seems specific to Gemma4.

And a huge THANK YOU for testing this. It is very much appreciated!!!

arthw · 2026-06-27T08:08:11Z

@johnkarlhill
There are several code to call wait().
Are they necessary to get the correct result?
Reduce or remove them will be quicker.

Thank you!

johnkarlhill · 2026-06-28T07:29:54Z

Bug fix

The MKL normalize kernel was writing output using a dense head-major layout (head * n_queries * DV), but llama.cpp's flash attention output uses an interleaved layout (query * n_heads + head per row, matching TILE's flash_attn_combine_results). Head 0 row 0 happened to alias at offset 0 in both layouts, so the first layer's first 64 output floats matched TILE. Everything else landed at wrong addresses. One-line fix in mkl_fa_normalize_head.

Tested models (all pass multi-turn coherence)

Gemma-4-26B-A4B-it (Q5_K_M, gqa=2)
Gemma-4-31B-it-qat (Q4_K_XL, dense)
Qwen3.6-27B (Q5_K_XL, gqa=6)
Qwen3.6-35B-A3B (Q4_K_XL, gqa=2)

Performance (Intel Arc Pro B70, Battlemage BMG-G21, 32K context, q8_0 KV cache)

Model	MKL PP (t/s)	TILE PP (t/s)	Speedup
Gemma-4-26B	1473	746	1.97×
Qwen3.6-27B	606	330	1.84×

Token generation unaffected (±1 t/s, within noise) — MKL only activates for prompt processing (n_kv ≥ 1024 with quantized KV).

Updated PR25025 - Qwen3.6-27B-MTP-UD-Q5_K_XL on B70.txt
Updated PR25025 - Gemma-4-26B-A4B-it-UD-Q5_K_M on B70.txt

How to test

cmake --preset x64-windows-sycl-release -DGGML_SYCL_F16=ON -DGGML_SYCL_DEVICE_ARCH=bmg_g21
cmake --build build-x64-windows-sycl-release --config Release -j 16

# Run (MKL activates automatically with flash-attn + quantized KV + n_kv ≥ 1024)
llama-server --flash-attn on --cache-type-k q8_0 --cache-type-v q8_0 --batch-size 1024 ...

# Disable for A/B comparison
set MKL_FA_DISABLE=1```

johnkarlhill · 2026-06-28T16:44:00Z

@johnkarlhill There are several code to call wait(). Are they necessary to get the correct result? Reduce or remove them will be quicker.

Thank you!

Removed 7 redundant stream->wait() calls — 3 no-ops covered by the SYCL in-order queue, 3 in a diagnostic block that was already gated behind MKL_FA_DIAG=1 (removed the whole block), and 1 final drain before pool destructors that's unnecessary with in-order semantics.

The 4 remaining waits are all required: oneMKL gemm() runs on its own internal queue that does not respect the SYCL in-order queue. Without these barriers, the softmax kernel would read stale KQ data and the accumulate kernel would read stale VKQ_chunk data. TILE uses zero explicit waits because everything is pure SYCL — MKL can't avoid these handshake points.

Perf unchanged. 609 (new) vs 606 t/s... within noise.

…essing

- Fix mkl_fa_normalize_head: use interleaved dst layout ((query * n_q_heads + head) * DV) matching TILE's flash_attn_combine_results. Previously used dense head-major layout which wrote head outputs to wrong addresses, corrupting attention for all models except Qwen3.6-27B (where GQA=6 heads were sparse enough to avoid visible overlap). - Remove 7 redundant stream->wait() calls — SYCL in-order queue already serializes pure SYCL kernel dependencies. Retain only the 4 MKL GEMM ↔ SYCL handshake barriers (oneMKL GEMM uses its own internal queue that does not respect SYCL in-order). - Remove unused dst_row_stride, diagnostic clutter, and dead K/V hex dump (fa_diag block in fattn-mkl.cpp). - Add MKL_FA_DISABLE=1 env var for A/B testing. - Add FA-DISP watchdog (MKL_FA_DEBUG=1) and FA-DIAG output fingerprint (MKL_FA_DIAG=1) in fattn.cpp. Tested: Gemma-4-26B, Gemma-4-31B, Qwen3.6-27B, Qwen3.6-35B-A3B Perf (B70/Battlemage, 32K, q8_0 KV): Gemma-4-26B: 1473 t/s MKL vs 746 TILE (1.97x) Qwen3.6-27B: 609 t/s MKL vs 330 TILE (1.85x) Co-Authored-By: Claude Code on DeepSeek-v4-Pro

arthw · 2026-06-29T01:53:32Z

+    static int fa_diag = -1;
+    static int fa_diag_count = 0;
+    if (fa_diag < 0) {
+        const char * e = getenv("MKL_FA_DIAG");


Rename:
MKL_FA_DISABLE to GGML_SYCL_ENABLE_MKL_FA.
MKL_FA_DIAG to GGML_SYCL_MKL_FA_DIAG

Explain them in SYCL.md, refer to chapter: # Environment Variable.

arthw · 2026-06-29T01:54:55Z

+            if (kb == BEST_FATTN_KERNEL_MKL) kname = "MKL";
+            if (kb == BEST_FATTN_KERNEL_TILE) kname = "TILE";
+            if (kb == BEST_FATTN_KERNEL_VEC) kname = "VEC";
+            fprintf(stderr, "[FA-DIAG] #%d %s D=%d n_kv=%lld n_q=%lld "


fprintf(stderr, ) is replaced by GGML_LOG_INFO()

arthw · 2026-06-29T01:55:52Z

+    // the same D — helps detect cache-truncation issues.
+    static int nkv_debug = -1;
+    if (nkv_debug < 0) {
+        const char * e = getenv("MKL_FA_DEBUG");


Rename:
MKL_FA_DEBUG to GGML_SYCL_MKL_FA_DEBUG
Explain it in SYCL.md

arthw · 2026-06-29T02:00:29Z

+        const char * e = getenv("MKL_FA_DISABLE");
+        mkl_disable = (e && e[0] == '1') ? 1 : 0;
+    }
+    if (mkl_disable == 0 && Q->ne[1] >= 128 && K->ne[1] >= 1024


if (mkl_disable == 0 && Q->ne[1] >= 128 && K->ne[1] >= 1024

Add comment to explain how to trigger FA_MKL in usage or llama-cli/server/bench parameters.

…, document in SYCL.md Completed the following: - Rename MKL_FA_DISABLE → GGML_SYCL_ENABLE_MKL_FA (inverted: 0 to disable) - Rename MKL_FA_DEBUG → GGML_SYCL_MKL_FA_DEBUG - Rename MKL_FA_DIAG → GGML_SYCL_MKL_FA_DIAG - Replace fprintf(stderr, ...) / fflush(stderr) with GGML_LOG_INFO() macro - Document all three env vars in docs/backend/SYCL.md under Runtime - Add comment explaining MKL FA activation trigger (flash-attn + quantized KV cache + batch-size >= 1024 + n_kv >= 1024) Resolves review feedback from arthw. Again, thank you!!! Co-Authored-By: Claude Code on DeepSeek-v4-Pro

arthw · 2026-06-29T05:55:28Z

+// MIT license
+// Copyright (C) 2025 Intel Corporation
+// SPDX-License-Identifier: MIT
+//


llama.cpp follow the unified copyright definition.
So, no need to declare here.

remove:

// // MIT license // Copyright (C) 2025 Intel Corporation // SPDX-License-Identifier: MIT //

arthw · 2026-06-29T05:59:51Z

 | GGML_SYCL_USE_LEVEL_ZERO_API | 1 (default) or 0 | Use Level Zero API for device memory allocation instead of SYCL. Reduces system RAM usage on Intel dGPUs by avoiding DMA-buf/TTM host memory staging. Requires GGML_SYCL_SUPPORT_LEVEL_ZERO_API=ON at build time. SYCL backend always runs on Level Zero running time even if it's set as OFF (The SYCL api will be usage for memory allocation).|
 | GGML_SYCL_DISABLE_DNN | 0 (default) or 1 | Disable running computations through oneDNN and always use oneMKL. |
 | GGML_SYCL_ENABLE_VMM | 0 or 1 (default) | Enable the virtual-memory device pool. |
+| GGML_SYCL_ENABLE_MKL_FA | 1 (default) or 0 | Enable oneMKL GEMM flash attention for XMX-accelerated prompt processing with quantized KV cache. Set to 0 to force the TILE kernel path for A/B testing. Activated at runtime when flash-attn is enabled (`-fa` or `--flash-attn on`), KV cache is quantized (`--cache-type-k/q *_0/*_1`), and KV length ≥ 1024. |


User can find the parameter: --cache-type-k in llama-cli.

But how to set for KV length ≥ 1024?
Please provide a detailed method in llama-cli/server for common user.

arthw · 2026-06-29T06:02:11Z

+#define MKL_TAKE_TIME(t0)  auto t0 = std::chrono::steady_clock::now()
+#define MKL_ACCUM(acc, t0) acc += (int64_t)std::chrono::duration_cast \


The two macros are used to debug for perf.
They should be disabled as default.

arthw · 2026-06-29T06:03:45Z

+                    stream->wait();
+                    try { ev.wait_and_throw(); } catch (sycl::exception & e) {


stream->wait() & ev.wait_and_throw() are duplicated wait() code.
Maybe impact side effect.
Remove one of them.

arthw · 2026-06-29T06:08:31Z

+                    stream->wait();
+                    try { ev.wait_and_throw(); } catch (sycl::exception & e) {


same comment as above: remove one.

arthw · 2026-06-29T06:10:16Z

+        const char * e = getenv("GGML_SYCL_ENABLE_MKL_FA");
+        mkl_disable = (e && e[0] == '0') ? 1 : 0;


The code can be replaced by existed function, like
ggml_sycl_get_env("GGML_SYCL_USM_SYSTEM", 0);

…ove dup waits, gate perf macros - Replace raw getenv() with ggml_sycl_get_env() in all 4 env-var checks (fattn.cpp: GGML_SYCL_ENABLE_MKL_FA, GGML_SYCL_MKL_FA_DEBUG, GGML_SYCL_MKL_FA_DIAG; fattn-mkl.cpp: GGML_SYCL_MKL_FA_DEBUG) - Remove duplicated stream->wait() before ev.wait_and_throw() in GEMM KQ and GEMM VKQ — ev.wait_and_throw() already waits for completion - Gate MKL_ACCUM macro behind do_print so timing accumulators are no-ops in normal operation - Remove redundant MIT/Intel copyright header from fattn-mkl.cpp - Remove unused #include <cfloat> - Expand SYCL.md MKL FA docs with step-by-step activation trigger and example llama-cli command Again, thank you!!! Co-Authored-By: Claude Code on DeepSeek-v4-Pro

Remove the quantized-only restriction on MKL activation — the MKL kernel converts any non-F16 K/V to F16 via to_fp16_sycl before GEMM, so F16 (default), BF16, and F32 caches all benefit from XMX hardware acceleration. The type restriction was an unnecessary gate. Before (F16/BF16 default cache + FA on at 32K prefill): ~356 t/s (TILE path) After: ~670 t/s (MKL path, matching quantized-cache baseline) Minimal change: two conditions removed, one comment updated in fattn.cpp. No kernel or conversion code changes — the dequant pipeline already covers all types.

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

…nv-var one-liners

…/llama.cpp into sycl-mkl-flash-attn

johnkarlhill · 2026-07-09T13:40:12Z

@arthw Updated based on your suggestions. Thank you!!

arthw

It's good job!

Thank you!

arthw · 2026-07-13T10:48:27Z

@johnkarlhill
Please resolve the conflict!
Thank you!

johnkarlhill · 2026-07-13T20:20:58Z

@arthw I think you meant the conflict with the new GGML_SYCL_ENABLE_FUSION env variable, so I added that. Please let me know if there was another conflict I needed to address.

NeoZhangJianyu · 2026-07-14T07:15:14Z

@johnkarlhill
The conflict means the github show conflict in the page.
It disappear now.

Thank you!

ggml-org#25025) * SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt processing * fattn-mkl: fix interleaved dst layout in normalize kernel - Fix mkl_fa_normalize_head: use interleaved dst layout ((query * n_q_heads + head) * DV) matching TILE's flash_attn_combine_results. Previously used dense head-major layout which wrote head outputs to wrong addresses, corrupting attention for all models except Qwen3.6-27B (where GQA=6 heads were sparse enough to avoid visible overlap). - Remove 7 redundant stream->wait() calls — SYCL in-order queue already serializes pure SYCL kernel dependencies. Retain only the 4 MKL GEMM ↔ SYCL handshake barriers (oneMKL GEMM uses its own internal queue that does not respect SYCL in-order). - Remove unused dst_row_stride, diagnostic clutter, and dead K/V hex dump (fa_diag block in fattn-mkl.cpp). - Add MKL_FA_DISABLE=1 env var for A/B testing. - Add FA-DISP watchdog (MKL_FA_DEBUG=1) and FA-DIAG output fingerprint (MKL_FA_DIAG=1) in fattn.cpp. Tested: Gemma-4-26B, Gemma-4-31B, Qwen3.6-27B, Qwen3.6-35B-A3B Perf (B70/Battlemage, 32K, q8_0 KV): Gemma-4-26B: 1473 t/s MKL vs 746 TILE (1.97x) Qwen3.6-27B: 609 t/s MKL vs 330 TILE (1.85x) Co-Authored-By: Claude Code on DeepSeek-v4-Pro * Thank you for the review feedback: rename env vars, use GGML_LOG_INFO, document in SYCL.md Completed the following: - Rename MKL_FA_DISABLE → GGML_SYCL_ENABLE_MKL_FA (inverted: 0 to disable) - Rename MKL_FA_DEBUG → GGML_SYCL_MKL_FA_DEBUG - Rename MKL_FA_DIAG → GGML_SYCL_MKL_FA_DIAG - Replace fprintf(stderr, ...) / fflush(stderr) with GGML_LOG_INFO() macro - Document all three env vars in docs/backend/SYCL.md under Runtime - Add comment explaining MKL FA activation trigger (flash-attn + quantized KV cache + batch-size >= 1024 + n_kv >= 1024) Resolves review feedback from arthw. Again, thank you!!! Co-Authored-By: Claude Code on DeepSeek-v4-Pro * Thank you for the review feedback round 2: use ggml_sycl_get_env, remove dup waits, gate perf macros - Replace raw getenv() with ggml_sycl_get_env() in all 4 env-var checks (fattn.cpp: GGML_SYCL_ENABLE_MKL_FA, GGML_SYCL_MKL_FA_DEBUG, GGML_SYCL_MKL_FA_DIAG; fattn-mkl.cpp: GGML_SYCL_MKL_FA_DEBUG) - Remove duplicated stream->wait() before ev.wait_and_throw() in GEMM KQ and GEMM VKQ — ev.wait_and_throw() already waits for completion - Gate MKL_ACCUM macro behind do_print so timing accumulators are no-ops in normal operation - Remove redundant MIT/Intel copyright header from fattn-mkl.cpp - Remove unused #include <cfloat> - Expand SYCL.md MKL FA docs with step-by-step activation trigger and example llama-cli command Again, thank you!!! Co-Authored-By: Claude Code on DeepSeek-v4-Pro * fattn-mkl: enable MKL FA for all KV cache types Remove the quantized-only restriction on MKL activation — the MKL kernel converts any non-F16 K/V to F16 via to_fp16_sycl before GEMM, so F16 (default), BF16, and F32 caches all benefit from XMX hardware acceleration. The type restriction was an unnecessary gate. Before (F16/BF16 default cache + FA on at 32K prefill): ~356 t/s (TILE path) After: ~670 t/s (MKL path, matching quantized-cache baseline) Minimal change: two conditions removed, one comment updated in fattn.cpp. No kernel or conversion code changes — the dequant pipeline already covers all types. * fattn-mkl: rename mkl_disable -> mkl_enable for clarity * fattn-mkl: refine MKL FA dispatch gates Three changes: 1. Remove quantized-only restriction - MKL FA activates for all KV cache types (F16 default, BF16, F32, quantized). The MKL kernel converts non-F16 K/V via to_fp16_sycl before GEMM. 2. Rename mkl_disable -> mkl_enable to match env var (GGML_SYCL_ENABLE_MKL_FA). 3. Replace batch-size threshold with Q->ne[1] >= 32 gate. Keeps TG (Q=1) and MTP drafts (Q=3-8) on VEC path where fused kernel beats MKL launch overhead. Routes all multi-token prefill through XMX-accelerated GEMM. Production data confirms Q patterns: 1-8 TG, 32-127 cache reuse, 128+ full reprocess. At 32K F16/BF16 FA-on: 356 -> 670 t/s. * ggml-sycl: fix F16 cache + MKL FA multi-turn corruption; add gate guards Two changes: 1. Always copy F16 K/V to dense row-major buffers before MKL GEMM. Previously F16 was read in-place with raw tensor strides. During multi-turn conversations, the accumulated KV cache had different stride properties than a fresh prefill, producing corrupted outputs. Now dense F16 gets a fast memcpy; interleaved (Gemma) gets a strided copy kernel. This matches what the quantized paths already did through to_fp16_sycl. 2. Gate MKL FA on unsupported op params (max_bias, logit_softcap, batch dim mismatch) and pathological F16 strides (nb[1] not a multiple of ne[0]*2). These conditions would previously crash inside the MKL kernel. Pathological strides (test-only) and ALiBi/softcap fall through to TILE/VEC which handle them correctly. The stride check uses modulo rather than equality, so both dense (nb1 == ne0*2) and interleaved (nb1 == H * ne0*2) pass — all real models use these layouts. Only test cases with overlapping rows (nb1=32 or nb1=75 for ne0=40) are blocked. Thanks to hmscider for the oneDNN FA PR (ggml-org#25222) which surfaced the same insight: always normalize inputs to contiguous F16 before GEMM. Co-Authored-By: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com> * fattn-mkl: fix quant+GQA KV strides, tighten MKL gate, add K>=1024 tests Adding K>=1024 flash-attn test cases surfaced several MKL bugs: - Quant K/V with a padded seq-view (real KV cache) used the wrong strides in the dequant path... only the true Gemma interleave layout should reconstruct strides. nb[2] vs ne[1]*nb[1] - Gate was firing on shapes the kernel doesn't handle: head_dim < 64 or not a multiple of 64, MHA, attention sinks, and bf16 decode... fell through to vec which no bf16 case. Gate MKL to the validated envelope: gqa>=2, head_dim 64 through 512 (has to be a multiple of 64) with matching K/V head size, mask, no sinks/alibi/softcap... everything else falls back to tile. Covers Qwen Dense/MoE and Gemma4 Dense/MoE Ran test-backend-ops -o FLASH_ATTN_EXT: 3641/3641 pass. Perplexity unchanged... 6.7267 MKL vs 6.7290 stock using Qwen 27b q5_k_xl * Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * fattn-mkl: bound attention scratch so it doesn't grow with batch or context... also dropped the bf16 comment in fattn.cpp per arthw review. * Update ggml/src/ggml-sycl/fattn-mkl.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn-mkl.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * apply arthw suggestions: enum for dequant modes, macro for wg_size, env-var one-liners --------- Co-authored-by: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com> Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

arthw · 2026-08-03T06:58:32Z

In code: commit 11924d4 (tag: b10223, origin/master, master)

It will get performance increase in more LLMs event with fp32 building on B60:

Test Script

'./build/bin/llama-bench --device SYCL0 -fa 1 -p 4096 -n 0 -m model.gguf'

Environment Configurations

env01
- export GGML_SYCL_ENABLE_MKL_FA=1
env02
- export GGML_SYCL_ENABLE_MKL_FA=0

Per Metric

model	metric	GGML_SYCL_ENABLE_MKL_FA=0	GGML_SYCL_ENABLE_MKL_FA=1
gemma-4-12b-it-Q5_K_M.gguf	pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1	172.73	372.13 (+115.44%)
qwen2-7b-instruct-q4_k_m.gguf	pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1	533.98	686.62 (+28.59%)
Qwen3-4B-Q4_K_M.gguf	pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1	527.50	617.69 (+17.10%)
Qwen3-8B-Q6_K.gguf	pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1	422.50	478.94 (+13.36%)
DeepSeek-R1-Distill-Llama-8B-Q4_0.gguf	pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1	458.13	517.34 (+12.92%)
Olmo-3-7B-Instruct-Q4_K_M.gguf	pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1	460.75	460.86 (+0.02%)
deepseek-moe-16b-chat.Q4_K_M.gguf	pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1	587.67	587.37 (-0.05%)
gpt-oss-20b-Q4_0.gguf	pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1	681.75	679.68 (-0.30%)
Bonsai-1.7B-Q1_0.gguf	pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1	1290.00	1024.29 (-20.60%)
granite-3.0-3b-a800m-instruct-Q4_0.gguf	pp4096, fa=1, dev=SYCL0, backend=SYCL, ngl=-1	1161.84	758.26 (-34.74%)

* test new flash_attn test * rebase and fix to disable subgrou matrices when max_kv_tile == 0 * delete log output * Add i32 support to cpy and enables the all ops test * restore the non target ci tests * comment out of TODO of build-cpu.yml * fix format

ggml-org#25025) * SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt processing * fattn-mkl: fix interleaved dst layout in normalize kernel - Fix mkl_fa_normalize_head: use interleaved dst layout ((query * n_q_heads + head) * DV) matching TILE's flash_attn_combine_results. Previously used dense head-major layout which wrote head outputs to wrong addresses, corrupting attention for all models except Qwen3.6-27B (where GQA=6 heads were sparse enough to avoid visible overlap). - Remove 7 redundant stream->wait() calls — SYCL in-order queue already serializes pure SYCL kernel dependencies. Retain only the 4 MKL GEMM ↔ SYCL handshake barriers (oneMKL GEMM uses its own internal queue that does not respect SYCL in-order). - Remove unused dst_row_stride, diagnostic clutter, and dead K/V hex dump (fa_diag block in fattn-mkl.cpp). - Add MKL_FA_DISABLE=1 env var for A/B testing. - Add FA-DISP watchdog (MKL_FA_DEBUG=1) and FA-DIAG output fingerprint (MKL_FA_DIAG=1) in fattn.cpp. Tested: Gemma-4-26B, Gemma-4-31B, Qwen3.6-27B, Qwen3.6-35B-A3B Perf (B70/Battlemage, 32K, q8_0 KV): Gemma-4-26B: 1473 t/s MKL vs 746 TILE (1.97x) Qwen3.6-27B: 609 t/s MKL vs 330 TILE (1.85x) Co-Authored-By: Claude Code on DeepSeek-v4-Pro * Thank you for the review feedback: rename env vars, use GGML_LOG_INFO, document in SYCL.md Completed the following: - Rename MKL_FA_DISABLE → GGML_SYCL_ENABLE_MKL_FA (inverted: 0 to disable) - Rename MKL_FA_DEBUG → GGML_SYCL_MKL_FA_DEBUG - Rename MKL_FA_DIAG → GGML_SYCL_MKL_FA_DIAG - Replace fprintf(stderr, ...) / fflush(stderr) with GGML_LOG_INFO() macro - Document all three env vars in docs/backend/SYCL.md under Runtime - Add comment explaining MKL FA activation trigger (flash-attn + quantized KV cache + batch-size >= 1024 + n_kv >= 1024) Resolves review feedback from arthw. Again, thank you!!! Co-Authored-By: Claude Code on DeepSeek-v4-Pro * Thank you for the review feedback round 2: use ggml_sycl_get_env, remove dup waits, gate perf macros - Replace raw getenv() with ggml_sycl_get_env() in all 4 env-var checks (fattn.cpp: GGML_SYCL_ENABLE_MKL_FA, GGML_SYCL_MKL_FA_DEBUG, GGML_SYCL_MKL_FA_DIAG; fattn-mkl.cpp: GGML_SYCL_MKL_FA_DEBUG) - Remove duplicated stream->wait() before ev.wait_and_throw() in GEMM KQ and GEMM VKQ — ev.wait_and_throw() already waits for completion - Gate MKL_ACCUM macro behind do_print so timing accumulators are no-ops in normal operation - Remove redundant MIT/Intel copyright header from fattn-mkl.cpp - Remove unused #include <cfloat> - Expand SYCL.md MKL FA docs with step-by-step activation trigger and example llama-cli command Again, thank you!!! Co-Authored-By: Claude Code on DeepSeek-v4-Pro * fattn-mkl: enable MKL FA for all KV cache types Remove the quantized-only restriction on MKL activation — the MKL kernel converts any non-F16 K/V to F16 via to_fp16_sycl before GEMM, so F16 (default), BF16, and F32 caches all benefit from XMX hardware acceleration. The type restriction was an unnecessary gate. Before (F16/BF16 default cache + FA on at 32K prefill): ~356 t/s (TILE path) After: ~670 t/s (MKL path, matching quantized-cache baseline) Minimal change: two conditions removed, one comment updated in fattn.cpp. No kernel or conversion code changes — the dequant pipeline already covers all types. * fattn-mkl: rename mkl_disable -> mkl_enable for clarity * fattn-mkl: refine MKL FA dispatch gates Three changes: 1. Remove quantized-only restriction - MKL FA activates for all KV cache types (F16 default, BF16, F32, quantized). The MKL kernel converts non-F16 K/V via to_fp16_sycl before GEMM. 2. Rename mkl_disable -> mkl_enable to match env var (GGML_SYCL_ENABLE_MKL_FA). 3. Replace batch-size threshold with Q->ne[1] >= 32 gate. Keeps TG (Q=1) and MTP drafts (Q=3-8) on VEC path where fused kernel beats MKL launch overhead. Routes all multi-token prefill through XMX-accelerated GEMM. Production data confirms Q patterns: 1-8 TG, 32-127 cache reuse, 128+ full reprocess. At 32K F16/BF16 FA-on: 356 -> 670 t/s. * ggml-sycl: fix F16 cache + MKL FA multi-turn corruption; add gate guards Two changes: 1. Always copy F16 K/V to dense row-major buffers before MKL GEMM. Previously F16 was read in-place with raw tensor strides. During multi-turn conversations, the accumulated KV cache had different stride properties than a fresh prefill, producing corrupted outputs. Now dense F16 gets a fast memcpy; interleaved (Gemma) gets a strided copy kernel. This matches what the quantized paths already did through to_fp16_sycl. 2. Gate MKL FA on unsupported op params (max_bias, logit_softcap, batch dim mismatch) and pathological F16 strides (nb[1] not a multiple of ne[0]*2). These conditions would previously crash inside the MKL kernel. Pathological strides (test-only) and ALiBi/softcap fall through to TILE/VEC which handle them correctly. The stride check uses modulo rather than equality, so both dense (nb1 == ne0*2) and interleaved (nb1 == H * ne0*2) pass — all real models use these layouts. Only test cases with overlapping rows (nb1=32 or nb1=75 for ne0=40) are blocked. Thanks to hmscider for the oneDNN FA PR (ggml-org#25222) which surfaced the same insight: always normalize inputs to contiguous F16 before GEMM. Co-Authored-By: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com> * fattn-mkl: fix quant+GQA KV strides, tighten MKL gate, add K>=1024 tests Adding K>=1024 flash-attn test cases surfaced several MKL bugs: - Quant K/V with a padded seq-view (real KV cache) used the wrong strides in the dequant path... only the true Gemma interleave layout should reconstruct strides. nb[2] vs ne[1]*nb[1] - Gate was firing on shapes the kernel doesn't handle: head_dim < 64 or not a multiple of 64, MHA, attention sinks, and bf16 decode... fell through to vec which no bf16 case. Gate MKL to the validated envelope: gqa>=2, head_dim 64 through 512 (has to be a multiple of 64) with matching K/V head size, mask, no sinks/alibi/softcap... everything else falls back to tile. Covers Qwen Dense/MoE and Gemma4 Dense/MoE Ran test-backend-ops -o FLASH_ATTN_EXT: 3641/3641 pass. Perplexity unchanged... 6.7267 MKL vs 6.7290 stock using Qwen 27b q5_k_xl * Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * fattn-mkl: bound attention scratch so it doesn't grow with batch or context... also dropped the bf16 comment in fattn.cpp per arthw review. * Update ggml/src/ggml-sycl/fattn-mkl.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn-mkl.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * apply arthw suggestions: enum for dequant modes, macro for wg_size, env-var one-liners --------- Co-authored-by: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com> Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

…ml-org#26566) * test new flash_attn test * rebase and fix to disable subgrou matrices when max_kv_tile == 0 * delete log output * Add i32 support to cpy and enables the all ops test * restore the non target ci tests * comment out of TODO of build-cpu.yml * fix format

43 upstream commits, 15 of them in our paths. Four conflicts, resolved as follows. The dangerous change in this range did NOT conflict -- see (2). 1. ggml/src/ggml-sycl/element_wise.cpp -- TOOK UPSTREAM VERBATIM. ggml-org#25946 landed upstream as 11b068d. We had been carrying it as a cherry-pick of the then-unmerged PR (ca7c42a) plus two commits of our own stacked on top: 27d821e fastdiv for the strided unary index reconstruction 0595878 fastdiv for the fused-GLU index reconstruction Upstream's landed form contains BOTH optimisations by the same mechanism (init_fastdiv_values host-side + fast_div_modulo in-kernel, on the strided unary path and on all five gated_op_fused_* kernels). All three of ours are therefore superseded and are dropped; the file is now byte-identical to origin/master. Only semantic difference we give up: ours guarded k > u32 with an exact int64 fallback, upstream asserts ggml_nelements(dst) < 2^31 instead -- stricter by 2x, and unreachable for a GLU activation (~8 GB at f32). Our ne>0?ne:1 divisor guard is also dropped; init_fastdiv_values asserts d != 0 and a ggml tensor always has ne[i] >= 1, so it was defensive, not load-bearing. 2. ggml/src/ggml-sycl/fattn.cpp -- PRECEDENCE PRESERVED, both kernels kept. Upstream ggml-org#25025 adds a oneMKL GEMM flash-attention path and gives it BEST_FATTN_KERNEL_MKL = 300 -- the value we already use for BEST_FATTN_KERNEL_MMA. git flagged the enum collision. It did NOT flag the consequential half: upstream places the MKL gate ABOVE our MMA check, and that hunk auto-merged clean. MKL's gate is default-ON (GGML_SYCL_ENABLE_MKL_FA=1) and its envelope -- gqa_ratio >= 2, head_dim % 64 in [64,512], Q->ne[1] >= 32, K->ne[1] >= 1024, no sinks / ALiBi / softcap, with a quantized K/V SKIPPING the F16 stride test -- matches our deploy prefill exactly. Taken verbatim it would have silently replaced the measured MMA kernel with an unmeasured one, staged the whole q8_0 KV cache to F16 first, and (per upstream's own note) broken SYCL graph capture replay. Resolution: MKL renumbered to 400 so both kernels stay reachable, and its gate takes an added !ggml_sycl_fattn_mma_supported(dst) conjunct. MMA wins where MMA is supported; MKL keeps its FULL envelope for everything MMA declines, which is upstream's intent in every case that is not ours. This is a precedence choice, not a revert -- and it is A/B-able without a rebuild: GGML_SYCL_FATTN_MMA=0 -> MMA declines, MKL takes the path GGML_SYCL_ENABLE_MKL_FA=0 -> MKL off entirely Also merged both sides' env-gated instruments, hoisting the kernel selection to a single call: upstream re-derived it three times (watchdog, switch, fingerprint), so an instrument could disagree with what actually ran. All three now read one hoisted `k`. Fixed a latent lie in our own FATTN_DEBUG printer while there -- BEST_FATTN_KERNEL_ONEDNN was printing as "NONE"; ONEDNN and MKL now print. 3. ggml/src/ggml-sycl/cpy.cpp -- kept ours. Upstream's side of the hunk was empty; ggml-org#26005 touched adjacent lines. Our GGML_SYCL_CPY_CENSUS instrument is unchanged. 4. tests/test-backend-ops.cpp -- kept both sides, additive and disjoint (same resolution as the 07-28 sync). Ours = the finding-85 deployed-shape MUL_MAT sweep; upstream's = m==1 either side of MMVF_MAX_BATCH_SIZE. Gates run before this commit: - cmake-option-audit.sh e9fa078 origin/master -> 7/7 watched options unchanged, and --selftest fires (rc=3), so the check is proven able to go red. - post-configure asserts: GGML_SYCL / _F16 / _DNN / _GRAPH all ON. NOT done in this commit, and required before any number from this tree is comparable to a pre-sync one: rebuild + re-baseline. Upstream changed code under every arm; ratios within one arm survive, absolutes do not.

…ml-org#26566) * test new flash_attn test * rebase and fix to disable subgrou matrices when max_kv_tile == 0 * delete log output * Add i32 support to cpy and enables the all ops test * restore the non target ci tests * comment out of TODO of build-cpu.yml * fix format

johnkarlhill requested a review from a team as a code owner June 26, 2026 01:24

github-actions Bot added ggml changes relating to the ggml tensor library for machine learning SYCL https://en.wikipedia.org/wiki/SYCL - GPU programming language labels Jun 26, 2026

arthw reviewed Jun 26, 2026

View reviewed changes

johnkarlhill force-pushed the sycl-mkl-flash-attn branch from a4871d8 to ce37155 Compare June 28, 2026 16:42

johnkarlhill added 2 commits June 28, 2026 18:23

SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt proc…

763c40f

…essing

arthw reviewed Jun 29, 2026

View reviewed changes

johnkarlhill force-pushed the sycl-mkl-flash-attn branch from ce37155 to 5a81c11 Compare June 29, 2026 03:10

github-actions Bot added the documentation Improvements or additions to documentation label Jun 29, 2026

arthw reviewed Jun 29, 2026

View reviewed changes

Comment thread ggml/src/ggml-sycl/fattn.cpp Outdated

Comment thread ggml/src/ggml-sycl/fattn.cpp Outdated

johnkarlhill added 2 commits June 29, 2026 10:12

johnkarlhill and others added 4 commits July 9, 2026 07:52

Update ggml/src/ggml-sycl/fattn-mkl.cpp

9ef4732

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

Update ggml/src/ggml-sycl/fattn-mkl.cpp

0808d31

Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>

apply arthw suggestions: enum for dequant modes, macro for wg_size, e…

75c6312

…nv-var one-liners

Merge branch 'sycl-mkl-flash-attn' of https://github.com/johnkarlhill…

b43b354

…/llama.cpp into sycl-mkl-flash-attn

arthw approved these changes Jul 10, 2026

View reviewed changes

Merge remote-tracking branch 'origin/master' into sycl-mkl-flash-attn

94d23ba

arthw added the merge ready A maintainer can use this label to indicate that they consider the changes final and ready to merge. label Jul 14, 2026

johnkarlhill mentioned this pull request Jul 17, 2026

SYCL: add HYBRID flash attention for quantized KV caches #25802

Closed

mndodd mentioned this pull request Jul 19, 2026

Extended SYCL oneDNN SDPA to non-FP16 KV caches (Q4_0–Q8_0 and FP32) #25874

Merged

1 task

ggerganov merged commit 9d9a6d2 into ggml-org:master Jul 31, 2026
26 of 28 checks passed

llin95 mentioned this pull request Aug 1, 2026

SYCL: b10216 oneMKL FA path regresses 200K prompt prefill on Arc Pro B70 by ~35% #26396

Closed

F3zz1k mentioned this pull request Aug 1, 2026

Eval bug: SYCL: oneMKL FA (#25025) slower with f16 KV, faster with q8_0 KV #26413

Closed

yomaytk mentioned this pull request Aug 4, 2026

ggml-webgpu: fix CI errors from #25025 and #25262 #26566

Merged

		#define MKL_TAKE_TIME(t0) auto t0 = std::chrono::steady_clock::now()
		#define MKL_ACCUM(acc, t0) acc += (int64_t)std::chrono::duration_cast \

		stream->wait();
		try { ev.wait_and_throw(); } catch (sycl::exception & e) {

		const char * e = getenv("GGML_SYCL_ENABLE_MKL_FA");
		mkl_disable = (e && e[0] == '0') ? 1 : 0;

Uh oh!

Conversation

johnkarlhill commented Jun 26, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

When it activates

Implementation

Testing

Known limitations

Debug output

AI disclosure

Uh oh!

ggml-gh-bot Bot commented Jun 26, 2026

Uh oh!

arthw left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

johnkarlhill commented Jun 26, 2026

Uh oh!

arthw commented Jun 26, 2026

Uh oh!

maxious commented Jun 26, 2026

Short context (pp512 — MKL path NOT active)

Long context (pp2048 — MKL path active)

Full commands

Uh oh!

johnkarlhill commented Jun 26, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

jlionhan commented Jun 26, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

johnkarlhill commented Jun 26, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

arthw commented Jun 27, 2026

Uh oh!

johnkarlhill commented Jun 28, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

johnkarlhill commented Jun 28, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

johnkarlhill commented Jul 9, 2026

Uh oh!

arthw left a comment

Choose a reason for hiding this comment

Uh oh!

arthw commented Jul 13, 2026

Uh oh!

johnkarlhill commented Jul 13, 2026

Uh oh!

NeoZhangJianyu commented Jul 14, 2026

Uh oh!

Uh oh!

arthw commented Aug 3, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Test Script

johnkarlhill commented Jun 26, 2026 •

edited

Loading

arthw left a comment •

edited

Loading

johnkarlhill commented Jun 26, 2026 •

edited

Loading

jlionhan commented Jun 26, 2026 •

edited

Loading

johnkarlhill commented Jun 26, 2026 •

edited

Loading

johnkarlhill commented Jun 28, 2026 •

edited

Loading

arthw commented Aug 3, 2026 •

edited

Loading