sycl: reduce redundant work in Q4_K multi-column MMVQ

[Eurekatic]

Overview

This fix was prompted by testing with DFlash, where performance with Q4 was slower than with Q8 and FP16. Upon investigating the cause, it became apparent that the reconstruction of the Q4_K weights—which was repeated in the destination columns in SYCL—could be optimized. This led me to examine how it was done in CUDA, where I observed a structural difference, including the use of two rows per block, which provided the clue to conduct an experimental verification.

Additional information

Kernel Performance

Q4_K N=3

upstream: 115.780 us
fix 1: 95.975 us
fix 1 + fix 2: 70.815 us

Fix 1: +17.11%
Fix 2: +26.22% incremental
Combined: +38.84%

Performance with Speculative Decoding

upstream: 42.2877 tok/s
fix 1: 49.9865 tok/s
fix 1 + fix 2: 54.9636 tok/s

Fix 1: +18.21%
Fix 2: +9.96% incremental
Combined: +29.98%

Correctness

Q4_K: 52/52 PASS
Q5_K: 27/27 PASS
fused Q4_K: 54/54 PASS

upstream, fix 1 and fix 1 + fix 2.
GRF: 128 -> 128 -> 128
spills: 0 -> 0 -> 0

Speculative Workload

684 cycles in all cases
1360 / 1368 accepted
acceptance = 99.415205%

Benchmark over 3 arms Identical byte-for-byte output across the three variants: upstream, fix 1 and fix 1 + fix 2.

Limitations of the Results

Specific to Q4_K: Enables blocking over 2-row by subgroup only for N=2..4; validation performed only over 2 B70s, I did not carry out cross-generation validation
Prefill: -0.19%

The significant improvement is evident in Speculative Decoding using a small-N multi-column path, whilst the generic tg128 benchmark only achieved an increase of approximately +0.46%.

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: YES
  I used AI for analysis, research and implementation; I identified the initial anomaly and defined the problem, leading to the design, execution and evaluation of the experiment. I verified the changes using A/B testing, correctness testing, output comparison and GPU resource checks. I have reviewed the changes and take responsibility for the submitted code.

[ggml-gh-bot]

Hi @Eurekatic, thanks for your contribution!

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

• PR Template not respected: Please respect the template when creating a new pull request. Make sure to fill out all required sections.

• AI-generated content: While code is allowed to be generated by AI, please write the PR description and commit messages on your own without the help of AI.

---

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

[WizardlyBump17]

Hello. I noticed a tg performance drop.

After 116k context:
5c1e50ffb8686a033ebbe661903993084589a728 (Pull Request) gives about 20t/s
a94d563 (your repo's master) gives about 27t/s

./llama-server --host 0.0.0.0 --port 8080 --model /models/Qwen3.6-35B-A3B-UD-Q5_K_XL.gguf.ignore --jinja --threads 8 --ctx-size 262144 --cache-ram 0 --parallel 1 --temperature 0.0 --top-p 0.2 --top-k 20 --load-mode none --spec-type draft-mtp --spec-draft-n-max 3 --batch-size 3650 --ubatch-size 3650 --n-gpu-layers 99 --n-cpu-moe 99 --ctx-checkpoints 0 --no-cache-prompt --cache-ram 0 --context-shift --cache-type-k q8_0 --cache-type-v q8_0

B580 + 7 5700X3D + 48GB at 3133MHz

[RaulAbejonDelgado]

Hello, thank you again for reporting this. We have been investigating the issue in considerable detail and have now managed to reproduce the performance regression using Qwen3.6-35B-A3B-UD-Q5_K_XL, the same quantization you used.

In our tests, with approximately 116k context, --n-cpu-moe 99 and MTP N=3, we get around 27.5 t/s on a94d563, which is very close to the ~27 t/s you reported. When testing the PR head, we also observe a performance drop, although on our hardware the magnitude is smaller than yours, at around 12%.

However, after testing the intermediate commits, we found something important: the regression does not appear to be introduced by the three Q4_K optimization commits in this PR.

a94d563 is actually the fork point of the PR branch. Between that commit and our optimization commits, the branch also picked up six upstream llama.cpp commits. When testing those states separately, the performance drop appears before our changes, at upstream commit 77918ca (server: allow accessing /metrics and /slots during llama_decode()).

With the same Q5_K_XL model and workload, we measured approximately:

a94d563 → 27.56 t/s
77918ca → 24.58 t/s
PR 5c1e50f → 24.19 t/s

With MTP disabled, we see practically no difference between these states. The regression appears specifically with the combination of MTP and CPU-resident MoE experts.

So your report did uncover a real performance regression, but our current results indicate that it is inherited from an upstream change between the PR fork point and our commits, rather than being introduced by the Q4_K optimizations themselves.

Thank you again for reporting this and for providing the complete configuration. It was very helpful in reproducing and isolating the issue. We are continuing to investigate the upstream change to see whether the regression can be fixed without losing the functionality it introduced.

[RaulAbejonDelgado]

We continued investigating the regression and were finally able to isolate quite precisely what was happening.

Our tests led us to the execution model introduced by commit 77918ca. In that state, llama_decode() is executed through the yield_to_queue worker thread. With the MoE experts resident on the CPU and MTP enabled, this change in execution context introduces a fairly clear additional cost on each cycle.

To verify that this was actually the cause, we ran a test keeping the same model, configuration, and workload, changing only the thread from which the decode work was executed. With the same Q5_K_XL model and approximately 116k context, the regressed state was around 161–165 ms/cycle, while keeping the compute on the original thread brought us back to approximately 144–145 ms/cycle, practically the same level we had before the regression appeared.

While investigating this, we also found that upstream had subsequently changed this mechanism in 22b8e31 (#27133, server: re-design yield_to_queue thread model). We decided to test that change directly and also compare it against its parent commit:

adb55e5 (parent) → 161.3 ms/cycle
22b8e31 (#27133) → 148.3 ms/cycle

This confirmed that the upstream redesign recovers practically all of the lost performance while, at the same time, keeping /metrics and /slots accessible during generation.

Once we had confirmed this, we updated this PR onto the latest upstream state that already contains #27133 and revalidated our three Q4_K optimization commits on top of this new base. In our tests, the previous regression no longer appears and the Q4_K optimizations continue to provide the expected improvements.

Since the performance loss you originally observed on your B580 was larger than on our hardware (~27 → ~20 t/s compared with ~27.5 → ~24 t/s in our initial tests), if you could test the current version of the PR again using the same configuration and workload you originally used, it would be very helpful to us.

This would also allow us to confirm on your B580 that the new PR base has restored the performance you had before.

Thx.

[WizardlyBump17]

We continued investigating the regression and were finally able to isolate quite precisely what was happening.

Our tests led us to the execution model introduced by commit 77918ca. In that state, llama_decode() is executed through the yield_to_queue worker thread. With the MoE experts resident on the CPU and MTP enabled, this change in execution context introduces a fairly clear additional cost on each cycle.

To verify that this was actually the cause, we ran a test keeping the same model, configuration, and workload, changing only the thread from which the decode work was executed. With the same Q5_K_XL model and approximately 116k context, the regressed state was around 161–165 ms/cycle, while keeping the compute on the original thread brought us back to approximately 144–145 ms/cycle, practically the same level we had before the regression appeared.

While investigating this, we also found that upstream had subsequently changed this mechanism in 22b8e31 (#27133, server: re-design yield_to_queue thread model). We decided to test that change directly and also compare it against its parent commit:

adb55e5 (parent) → 161.3 ms/cycle 22b8e31 (#27133) → 148.3 ms/cycle

This confirmed that the upstream redesign recovers practically all of the lost performance while, at the same time, keeping /metrics and /slots accessible during generation.

Once we had confirmed this, we updated this PR onto the latest upstream state that already contains #27133 and revalidated our three Q4_K optimization commits on top of this new base. In our tests, the previous regression no longer appears and the Q4_K optimizations continue to provide the expected improvements.

Since the performance loss you originally observed on your B580 was larger than on our hardware (~27 → ~20 t/s compared with ~27.5 → ~24 t/s in our initial tests), if you could test the current version of the PR again using the same configuration and workload you originally used, it would be very helpful to us.

This would also allow us to confirm on your B580 that the new PR base has restored the performance you had before.

Thx.

Seems to be fixed. Thank you

[arthw]

It's good job!

With Qwen3.5-4B-Q4_K_M.gguf on B60:

Test	fa	Base t/s	Primary t/s	Increase Rate (Primary vs Base)
pp512	0	267.02	302.27	13.20%
pp512	1	265.32	300.04	13.09%
tg128	0	88.05	88.71	0.75%
tg128	1	88.38	88.93	0.62%

Thank you!

[Titaniumtown]

@Eurekatic Thank you for the PR! Some comments:

1. Could you add a Q4_K case for perf -o MUL_MAT so this can be more easily tested?

2. A lot of references to the value of 6272, what is this value? It is not explained anywhere how it was derived, and if we are to have a variable like this, it should be #define'd somewhere with a comment explaining how it was derived.

2a. If we are to have a magic value like this, we should have a test for above and below nrows of 6272 to ensure that both code paths work.

I'm going to read over it again later.

[Titaniumtown]

This could be changed to be similar to: https://github.com/ggml-org/llama.cpp/pull/27062/changes#diff-234937ee14469de2d98d2bd6f0a382992f291c322090aa86705f819ecb3f3d62R124

maybe:

const auto a = reorder_vec_dot_q_sycl::load_activations(q8_1_quant_ptr, q8_1_ds_ptr, iqs);
partial_sum[j][0] += reorder_vec_dot_q_sycl::apply(wx, a);
partial_gate[j][0] += reorder_vec_dot_q_sycl::apply(wg, a);

[RaulAbejonDelgado]

Hello! Thank you for your questions.
in stages:
1-Yes it's possible,we have added Q4_K coverage to perf of MUL_MAT.
2-A:
We added constexpr int64_t Q4_K_MMVQ_ROW_PAIR_MIN_NROWS = 6272;
ORIGINAL SWEEP
6144 → OLD
6272 → ROW
crossover ∈ (6144, 6272]

FOLLOW-UP SWEEP
6144 → OLD
6176 → ROW
refined crossover ∈ (6144, 6176]

DECISION
keep 6272 as a conservative gate

We keep 6272. Why ? Because 6272 clearly favours ROW, we decided to adopt a conservative approach; we did not narrow the gate during this follow-up; 6176–6271 remains a possible area for further optimization

Q4_K MUL_MAT
ncols = 17408
ncols_dst = 2
Arc Pro B70
median µs/run

| nrows | OLD µs | ROW µs | ROW delta | Winner |
| 6144 | 136.12 | 149.42 | +9.77% | OLD |
| 6176 | 165.27 | 150.05 | −9.21% | ROW |
| 6208 | 153.38 | 136.66 | −10.91% | ROW |
| 6272 | 158.46 | 147.68 | −6.80% | ROW |
| 6400 | 170.03 | 149.08 | −12.32% | ROW |

2a.
We added boundary coverage around 6272 (6271, 6272, 6273) to check both sides of the gate, in both MUL_MAT and MUL_MAT_VEC_FUSION

load_activations()/apply():

That’s correct – the activation is shared, but it is applied independently to the two weights (up and gate). Now the activation is loaded just once; after that, its reuse is independent, and the mathematical calculation is not altered at any point.

[Titaniumtown]

@RaulAbejonDelgado can you rebase? Thank you!

[Eurekatic]

@Titaniumtown Rebased onto the latest master. I also rebuilt the SYCL backend and re-tested the Q4_K path on Intel Arc Pro B70 after the rebase. The focused correctness tests around the 6272-row threshold pass (6271/6272/6273), and a Q4_K_M inference smoke test also completed successfully. Thank you!

[arthw]

@Eurekatic
Could you change the status from Draft to Open if you think it's ready to review and merge?

Thank you!

[Titaniumtown]

LGTM

[ggerganov]

These tests are failing randomly because the max NMSE error is not correctly defined. You can easily reproduce with this patch:

diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 154146dda..954ac5b41 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -10383,8 +10383,12 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     for (ggml_glu_op glu_op : { GGML_GLU_OP_SWIGLU, GGML_GLU_OP_GEGLU }) {
         for (int64_t m_batch : { 2, 3, 4 }) {
             for (int64_t rows : { 1, 2, 3 }) {
+                if (rows == 1) {
+                    for (int p = 0; p < 1000; ++p) {
                 test_cases.emplace_back(new test_mul_mat_vec_fusion(GGML_TYPE_Q4_K, glu_op, m_batch, rows, 256,
                     false, 16, 8, false, false, true, false, { 1, 1 }));
+                    }
+                }
             }
         }
     }

Try to fix this. In the meantime, I will disable them in #28306

Also be mindful when adding new tests - it's good to be exhaustive, but the runtime keeps increasing with more and more tests.

[RaulAbejonDelgado]

Thanks for the reproducer and for pointing this out. I'll investigate the NMSE threshold for these cases and work on a minimal fix so the tests can be re-enabled.

I will also keep the test runtime in mind for future coverage additions. Thanks!

[RaulAbejonDelgado]

Thanks for the reproducer. I've spent some more time looking into this and I can reproduce the intermittent failures in the small Q4_K + GLU cases on SYCL (Intel Arc Pro B70).

First I compared the same tests before and after #27062:

                    pre-#27062    post-#27062

Runs 6000 6000
Failures 188 180
Failure rate 3.13% 3.00%

So at least statistically I don't see a regression introduced by #27062.

I wanted to make sure this wasn't just hidden by the random inputs, so I temporarily added a deterministic seed (GGML_TEST_DETERMINISTIC_SEED) and looked for cases I could reproduce exactly.

I found two GEGLU, rows=1, m_batch=3 outliers with seeds 36 and 37. Running those exact inputs before and after #27062 gives identical inputs, reference tensors and SYCL outputs.

seed 36:
reference_energy = 0.621023583
squared_error = 0.009189383
RMS error = 0.055345529
NMSE = 0.014797156

seed 37:
reference_energy = 1.216148594
squared_error = 0.016393911
RMS error = 0.073923183
NMSE = 0.013480187

Both are above the current max_nmse_err() = 5e-3, but importantly they behave exactly the same pre/post #27062.

I also checked the cases around Q4_K_MMVQ_ROW_PAIR_MIN_NROWS = 6272 again:

rows=6271 -> PASS
rows=6272 -> PASS
rows=6273 -> PASS

My suspicion at this point is that the fixed 5e-3 NMSE limit isn't a great fit for these very small fused GLU cases. When the reference energy gets small, NMSE can become quite sensitive even though the absolute error isn't proportionally large.

I experimented with an absolute + relative RMS comparison as well:

RMS(error) <= abs_tol + rel_tol * RMS(reference)

It can be tuned to accept the deterministic outliers above, but I wasn't comfortable stopping there. I injected known errors into captured outputs to see what we would lose by using those tolerances.

The result wasn't good enough IMO. The tolerances needed to accept the outliers also made the test noticeably less sensitive to additive and localized errors. In one of the tests it still accepted localized corruption affecting roughly 25% of the output. Sign inversion was detected, but overall I don't think this is safe enough to replace the current NMSE check.

So I don't think simply relaxing max_nmse_err() is the right fix either.

One thing that does look safe is reducing the small test matrix from 18 cases to 6:

GLU = { SWIGLU, GEGLU }
rows = { 1, 2, 3 }
m_batch = 3

and keeping the 6271 / 6272 / 6273 boundary cases.

That keeps one-row/even/odd small shapes for both GLU ops without testing the full m_batch={2,3,4} product.

So where I've ended up is: the flaky small-row behaviour is already present before #27062, the deterministic cases I've found are identical before/after the change, and the 6271/6272/6273 boundary tests still pass. There does seem to be an issue with how NMSE behaves for some of these small outputs, but the alternative I've tested so far weakens error detection too much, so I'd rather not propose a tolerance just to make the tests pass.

If there's an error metric or an existing test pattern you'd prefer for this kind of Q4_K fused output, I'm happy to implement it and test it.