Skip to content

metal : few-row MMA mat-mul and batched copies for speculative decoding - #29869

Merged
ggerganov merged 10 commits into
ggml-org:masterfrom
pratiknarola-t:metal-few-row-mma
Oct 5, 2026
Merged

ggerganov merged 10 commits into
ggml-org:masterfrom
pratiknarola-t:metal-few-row-mma

Conversation

@pratiknarola-t

@pratiknarola-t pratiknarola-t commented Oct 2, 2026 •

Copy link
Copy Markdown
Contributor

Overview

Speculative decoding verifies a few draft tokens per step, and batched decoding runs one row per sequence, so the model runs mat-muls with 2..16 src1 rows. On Apple GPUs without the tensor API (M1 to M4), Metal runs these with the mat-vec kernels, whose time grows with each row. On an M3 Ultra, DFlash2 decoding of Qwen3.8-27B is therefore slower than serial decoding on master.

This PR adds:

  • Mat-mul kernels for 2..16 src1 rows on 8x8 simdgroup matrices. Each weight is dequantized once for all rows, and the simdgroups of a threadgroup split K. Q4_0, Q8_0 and Q5_K have their own kernels. F32, F16, Q4_1, Q5_0, Q5_1, Q4_K and Q6_K use a generic path over the existing 16-weight dequantizers. Q4_0 at 2 rows uses a 2-row variant of the mat-vec kernel.
  • The kernels fill simdgroup matrices per lane, so they run only on MTLGPUFamilyApple7+ without the tensor API. They start at the row count where they beat the mat-vec kernels on an M3 Ultra: 6 for F32, 3 for F16, Q4_K, Q5_0 and Q5_1, and 2 for the other types.
  • A fusion table entry: MUL_MAT + ADD writes the residual sum from the MMA store.
  • CONCAT splits long rows across threadgroups when there are few rows.

Additional information

Changes to the fusion and reorder code:

  • The fusion checks, ggml_metal_fusion_next, ggml_metal_fusion_max and ggml_graph_optimize take the device props, so the reorder packs MUL_MAT + ADD only on devices that fuse it. The pack does not depend on the src1 row count, so graphs of different batch sizes get the same node order (no reallocation with GGML_SCHED_NO_REALLOC).
  • The alloc-deps pass also matches the new MUL_MAT + ADD pattern, so on every Metal device the mat-mul inputs stay alive until the ADD.

Performance on an Apple M3 Ultra (60-core GPU), macOS 15.7.9, Qwen3.8-27B Q4_0, master 836d571.

test-backend-ops perf -o MUL_MAT, m=4096, k=14336, time per run of this PR divided by master (two interleaved runs each). This table is from master 2923cf2. The kernels of this PR and the mat-vec kernels of master did not change after that.

type n=2 n=3 n=4 n=5 n=8
F32 1.00 0.97 1.00 0.99 0.57
F16 0.99 0.97 0.71 0.59 0.34
Q4_0 0.58 0.58 0.44 0.36 0.22
Q4_1 0.93 0.65 0.50 0.41 0.25
Q5_0 1.01 0.95 0.75 0.63 0.38
Q5_1 0.99 0.94 0.75 0.63 0.37
Q8_0 0.86 0.59 0.46 0.38 0.24
Q4_K 1.00 0.97 0.56 0.45 0.29
Q5_K 0.81 0.56 0.49 0.45 0.25
Q6_K 0.92 0.64 0.48 0.39 0.21

n=1 and n=512 are unchanged (0.99 to 1.01). A sweep of 9..16 rows gives 0.25 to 0.40 for all ten types.

llama-batched-bench -npp 512 -ntg 32 -npl 1,2,3,4,8 -c 32768 -pps -kvu, TG t/s (mean of two interleaved runs each). PP is 317 t/s on both builds.

B master this PR
1 32.9 32.9
2 32.7 49.7
3 36.2 62.2
4 36.3 76.1
8 38.6 119.7

llama-server with a Qwen3.8-27B DFlash2 Q8_0 drafter, -ngl 99 -fa on -c 8192 -np 1 --jinja, DFlash2 with --spec-type draft-dflash --spec-draft-n-max 7. 64 generated tokens, median of 5 requests after one warm-up, mean of two server runs. Decode tok/s:

mode prompt T master this PR
serial code 0 32.1 32.0
serial code 1 32.1 32.0
serial prose 0 32.1 32.0
serial prose 1 32.1 32.0
DFlash2 code 0 30.2 110.0
DFlash2 code 1 24.3 80.9
DFlash2 prose 0 16.8 62.6
DFlash2 prose 1 13.9 48.8

At T=0 the DFlash2 replies are the same as the serial replies (first 160 characters). At T=1 the two builds accept the same draft tokens.

llama-bench -fa 1 gives pp512 317.2 t/s on both builds, and tg128 32.45 t/s on master and 32.36 t/s with this PR (two interleaved runs each). llama-perplexity (10 chunks of 512 tokens) gives 5.1775 on both builds with -ub 512, and 5.1773 on master and 5.1774 with this PR with -ub 8.

Tests:

  • test-backend-ops -b MTL0: 16806/16806 on the M3 Ultra, and on an Apple M5 with the tensor API on and off (GGML_METAL_TENSOR_DISABLE=1).
  • test-metal-graph-optimize passes on both devices.
  • With -DGGML_SCHED_NO_REALLOC=ON, as in the Metal CI job: test-llama-archs -s 1 with 1 to 4 devices gives the same results as master, with no failures, and test-fusion --check tests/fusion/MTL.csv passes. The test-fusion graphs have 1 or 32 src1 rows, so the new fusion does not fire there.

The row thresholds come from the M3 Ultra. With the tensor API turned off, an M5 reaches parity one or two rows later for some types. I did not tune for that, because the M5 uses the tensor API by default.

This ports the Metal part of tetherto#321 to master.

Requirements

  • I have read and agree with the contributing guidelines
  • AI usage disclosure: YES - the kernels, the port to master, the tests, the benchmarks and this description were written with Claude Code (Anthropic).

Speculative decoding verifies a few draft tokens per step. Without the tensor API, Metal ran these mat-muls with the mat-vec kernels, whose time grows with every src1 row, so DFlash2 decoding on an M3 Ultra was slower than serial decoding.

- add mat-mul kernels for 2..16 src1 rows on 8x8 simdgroup matrices: each weight is dequantized once for all rows, and the simdgroups of a threadgroup split K. Q4_0, Q8_0 and Q5_K have their own kernels, F32, F16, Q4_1, Q5_0, Q5_1, Q4_K and Q6_K use a generic path over the 16-weight dequantizers, and Q4_0 at 2 rows uses a 2-row variant of the mat-vec kernel
- use them only on MTLGPUFamilyApple7+ without the tensor API, from the row count at which they beat the mat-vec kernels on an M3 Ultra (F32: 6, F16, Q4_K, Q5_0, Q5_1: 3, other types: 2)
- fusion table: MUL_MAT + ADD adds a same-shape residual in the MMA store, and up to 16 adjacent same-layout f32 copies between the same two tensors run as one dispatch
- the fusion checks and ggml_graph_optimize take the device props, so the reorder packs MUL_MAT + ADD only on devices that can fuse it, at every src1 row count
- views do not count toward GGML_METAL_FUSION_MAX when the reorder packs a group, so 16 recurrent state snapshot copies with views between them stay one group
- the encoder checks the inner nodes of a fused group for concurrency, tracks written views by their extent, and does not count the destination of a CPY as a read
- CONCAT splits long rows across threadgroups when there are few rows
- tests: few-row MUL_MAT, MUL_MAT_ADD, CPY_BATCH and CONCAT cases in test-backend-ops (with a prepare_graph hook for the copy order), test-metal-graph-optimize, test-metal-cpy-batch-alias
@github-actions github-actions Bot added testing Everything test related ggml changes relating to the ggml tensor library for machine learning Apple Metal https://en.wikipedia.org/wiki/Metal_(API) labels Oct 2, 2026
@ggml-gh-bot

ggml-gh-bot Bot commented Oct 2, 2026

Copy link
Copy Markdown

Hi @pratiknarola-t, thanks for your contribution!

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

  • Large PR: Large changes require prior discussion (e.g. an issue or RFC) and maintainers may not be able to review this PR as-is. Consider splitting it into smaller, focused PRs.

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

@ggerganov ggerganov self-assigned this Oct 2, 2026
Comment on lines +853 to +854
// longest batch first, so ggml_metal_fusion_next checks no shorter batch once one matches
{ GGML_METAL_FUSION_CPY_BATCH, std::vector<ggml_op>(16, GGML_OP_CPY), {}, true, ggml_metal_fusion_check_cpy_batch },

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do you have an estimate of how big is the impact of the CPY_BATCH fusion logic on the performance? The logic is a bit complicated and I'm wondering if it is worth having it.

@ggerganov ggerganov Oct 2, 2026 •

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

On M2 Ultra I don't measure noticeable improvement from the CPU_BATCH fusion:

llama-batched-bench -hf ggml-org/Qwen3.8-27B-GGUF:Q4_K_M -npp 512 -ntg 32 -npl 1,2,3,4,8 -c 32768 -pps -kvu

PR:

PP TG B N_KV T_PP s S_PP t/s T_TG s S_TG t/s T s S t/s
512 32 1 544 1.454 352.04 1.226 26.10 2.680 202.96
512 32 2 576 1.450 353.02 1.963 32.61 3.413 168.77
512 32 3 608 1.452 352.54 2.131 45.06 3.583 169.69
512 32 4 640 1.450 353.06 2.253 56.82 3.703 172.84
512 32 8 768 1.451 352.93 2.734 93.63 4.185 183.51

PR + CPY_BATCH fusion removed:

PP TG B N_KV T_PP s S_PP t/s T_TG s S_TG t/s T s S t/s
512 32 1 544 1.459 351.00 1.234 25.92 2.693 202.00
512 32 2 576 1.456 351.64 1.978 32.36 3.434 167.75
512 32 3 608 1.455 351.89 2.144 44.77 3.599 168.93
512 32 4 640 1.454 352.08 2.262 56.59 3.716 172.21
512 32 8 768 1.450 353.12 2.742 93.35 4.192 183.19
Patch:
diff --git a/ggml/src/ggml-metal/ggml-metal-fusion.cpp b/ggml/src/ggml-metal/ggml-metal-fusion.cpp
index 2f801b8498..72190e53f4 100644
--- a/ggml/src/ggml-metal/ggml-metal-fusion.cpp
+++ b/ggml/src/ggml-metal/ggml-metal-fusion.cpp
@@ -851,21 +851,21 @@ static const std::vector<ggml_metal_fusion> ggml_metal_fusions = {
     { GGML_METAL_FUSION_SSM_CONV_SILU,  ops_ssm_conv_silu,          {},     false, ggml_metal_fusion_check_ssm_conv_silu },
     { GGML_METAL_FUSION_MUL_MAT_ADD,   ops_mul_mat_add,            {},     false, ggml_metal_fusion_check_mul_mat_add },
     // longest batch first, so ggml_metal_fusion_next checks no shorter batch once one matches
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>(16, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>(15, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>(14, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>(13, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>(12, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>(11, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>(10, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 9, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 8, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 7, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 6, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 5, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 4, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 3, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
-    { GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 2, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>(16, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>(15, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>(14, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>(13, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>(12, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>(11, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>(10, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 9, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 8, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 7, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 6, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 5, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 4, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 3, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
+    //{ GGML_METAL_FUSION_CPY_BATCH,     std::vector<ggml_op>( 2, GGML_OP_CPY), {}, true,  ggml_metal_fusion_check_cpy_batch },
 };
 
 static_assert(GGML_METAL_CPY_BATCH_MAX <= GGML_METAL_FUSION_MAX, "a copy batch must fit in one fusion");

So I would recommend to simplify by removing it together with the memory ranges changes.

Remove the batched copy fusion with its kernel and tests, and revert the
memory range changes, as suggested in review. The memory ranges, the
graph reorder and the CPY encoder are again the same as on master.
# Conflicts:
#	ggml/src/ggml-metal/ggml-metal-impl.h
Comment on lines +659 to +661
// MUL_MAT + ADD of an f32 non-weight: the reorder packs it without reading row counts, so ubatch sizes share one order;
// the encoder fuses only a same-shape residual in the few-row MMA store, which the sum may overlap only in place
static bool ggml_metal_fusion_check_mul_mat_add(

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What is the reason to fuse the ADD only in the MMA case?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The ADD fusion was written together with the new few-row kernels because it was nearly free there. The MMA kernel already ends in one store loop, after the threadgroup reduction of its K slices. The fused residual is one function constant and one extra load per output element there (mul_mv_mma.metal:66-108). The PR targets the 2..16-row path, so only that store got the epilogue. Nothing about the ADD fusion is MMA-specific.

Comment thread ggml/src/ggml-metal/ggml-metal-common.h Outdated
Comment thread tests/test-metal-graph-optimize.cpp
Comment thread ggml/src/ggml-metal/ggml-metal-fusion.cpp
@ggerganov

Copy link
Copy Markdown
Member

@pratiknarola-t I'll push a change to this branch if that is OK - would like to do some clean-up, naming, etc.

@ggerganov

Copy link
Copy Markdown
Member

The row thresholds come from the M3 Ultra. With the tensor API turned off, an M5 reaches parity one or two rows later for some types. I did not tune for that, because the M5 uses the tensor API by default.

Based on my tests, enabling the MMA path on M5 devices has significant benefit - similar to the the reported benefit for M3 Ultra. It makes sense because even when the tensor API is available, we still use mul_mv_ext on master which we see is inferior to the mul_mv_mma for small batch sizes. So I will also update the logic to enable MMA for M5 too.

@ggerganov

Copy link
Copy Markdown
Member

I did some clean-up and simplified the logic:

  • Moved the new MMA kernels in new mul_mv_mma.metal source file
  • Simplified the fusion by removing the device properties. The idea is during graph optimize, we optimistically assume fusion is possible regardless of hardware constraints (for example, it does not matter if the device supports simdgroup_mm). This is much simpler and has negligible impact on perf for old generation hardware.
  • Enabled mul_mv_mma even when tensors cores are available (i.e. M5+)
  • Removed the special-case for Q4_0 with 2 src1 rows - this is too much logic for a very narrow case. We can reconsider a cleaner/generic implementation in a follow-up PR
  • Various renaming and code movement/refactor

Unless I accidentally broke something, I think this PR is now ready to be merged. The performance improvement for parallel generation and speculative decoding is significantly improved, so I would like to get this merged soon to make it for the v0.6.0 release tomorrow.

@ggerganov
ggerganov marked this pull request as ready for review October 4, 2026 11:21
@ggerganov
ggerganov requested a review from a team as a code owner October 4, 2026 11:21
@pratiknarola-t

Copy link
Copy Markdown
Contributor Author

Thanks a lot @ggerganov

@ggerganov
ggerganov merged commit a3a1c47 into ggml-org:master Oct 5, 2026
37 of 40 checks passed
@ggerganov

ggerganov commented Oct 5, 2026 •

Copy link
Copy Markdown
Member

@pratiknarola-t Thanks. Do you plan writing similar kernels for MUL_MAT_ID? And expanding the type coverage for the MMA kernels here?

@ggerganov ggerganov added the highlight Changes that will be highlighted in the next release notes label Oct 5, 2026
@pratiknarola-t

Copy link
Copy Markdown
Contributor Author

Yes @ggerganov. I am currently working on some optimisations. Once its completed, I will pick it up for MUL_MAT_ID as well.

ravi9 added a commit to ravi9/llama.cpp that referenced this pull request Oct 5, 2026
Op tests build Q4_1/Q4_K weights as u4 with an f16 zero point. The GPU
plugin fails to compile that form for some row counts with "clFinish,
error code: -5 CL_OUT_OF_RESOURCES", which aborts test-backend-ops on
the MUL_MAT cases added in ggml-org#29869 (e.g. m=1000, n=2, k=1024). Model
weights use a u4 zero point and are not affected.

Report these cases as unsupported on GPU until the plugin is fixed.
Op tests check support before allocating, so the check matches unbound
weights only; model loading probes with a dummy buffer and keeps its
weights on the GPU.
ravi9 added a commit to ravi9/llama.cpp that referenced this pull request Oct 6, 2026
Op tests build Q4_1/Q4_K weights as u4 with an f16 zero point. The GPU
plugin fails to compile that form for some row counts with "clFinish,
error code: -5 CL_OUT_OF_RESOURCES", which aborts test-backend-ops on
the MUL_MAT cases added in ggml-org#29869 (e.g. m=1000, n=2, k=1024). Model
weights use a u4 zero point and are not affected.

Report these cases as unsupported on GPU until the plugin is fixed.
Op tests check support before allocating, so the check matches unbound
weights only; model loading probes with a dummy buffer and keeps its
weights on the GPU.
ggerganov pushed a commit that referenced this pull request Oct 6, 2026
* ggml-openvino: skip unselected graph branches and support DUP

Upstream #29622 adds a mixed token/embd branch to every input
embedding graph through ggml_build_forward_select(). Its nodes are
not flagged for compute, but the backend translated them anyway,
and the DUP in that branch was unsupported, so the scheduler split
the graph and passed the embeddings across the split with a fixed
token count. The first single-token decode then failed
(test-thread-safety on CPU and GPU).

Build the OV model from the compute nodes only, and translate a
same-type contiguous DUP like CONT so the graph stays on one backend.

* ggml-openvino: make inp_scale_rows token dim dynamic

#29622 also moves the per-token embedding scale (gemma3, gemma3n,
gemma4) into a new [1, n_tokens] input. Give it a dynamic token dim
and pad it per chunk on the static (NPU) path.

* ggml-openvino: skip GPU MUL_MAT op tests with unbound Q4_1/Q4_K weights

Op tests build Q4_1/Q4_K weights as u4 with an f16 zero point. The GPU
plugin fails to compile that form for some row counts with "clFinish,
error code: -5 CL_OUT_OF_RESOURCES", which aborts test-backend-ops on
the MUL_MAT cases added in #29869 (e.g. m=1000, n=2, k=1024). Model
weights use a u4 zero point and are not affected.

Report these cases as unsupported on GPU until the plugin is fixed.
Op tests check support before allocating, so the check matches unbound
weights only; model loading probes with a dummy buffer and keeps its
weights on the GPU.

* ggml-openvino: create FILL in the output type

translate_fill always built an f32 constant, so an f16 FILL produced
f32 data and the copy back overran the f16 output buffer. Use the
output type for the constant.

* ggml-openvino: reject CONCAT with a quantized type

Quantized inputs are dequantized when translated, so the backend cannot
write a quantized CONCAT output. Report it as unsupported, as for CPY
to a quantized type.

* ggml-openvino: handle the single recurrent state gather of build_rs

#29856 changed build_rs to gather all recurrent states with one GET_ROWS
on the s_copy leaf and take the ubatch and extra states as views of it.
The stateful path matched only the previous form, a GET_ROWS per view of
s_copy, so Qwen3.5 failed with stateful execution on CPU and GPU
("is_axis_valid(axis, r)" in a Concat).

For a single-slot cache, treat the GET_ROWS on the s_copy leaf as the
active-state gather, keep the rank-4 layout of reshapes that read a view
of it, and map the copy of the empty extra-state view to the single-slot
remainder writeback. Do not warn about the dynamic dim of empty views.

* openvino: align eltwise operand ranks to work around a GPU-plugin defect

* openvino: match the MoE fusion on the rank-3 stateful graph

* ggml-openvino: do not unsqueeze an RMS norm output in AlignEltwiseOperandRanks

The pass unsqueezes the lower-rank operand of an Add/Multiply/Subtract
whose operand ranks differ. In gemma-3 the lower-rank operand of the
post-attention residual add is the norm output, and unsqueezing it makes
the GPU plugin compute the layer wrongly: gemma-3 returns empty answers
on GPU with stateful execution.

Skip the rewrite when the lower-rank operand is an RMS norm output.

* docs : update OpenVINO validated models

---------

Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com>
edwardyoon pushed a commit to edwardyoon/focus-llama that referenced this pull request Oct 8, 2026
* ggml-openvino: skip unselected graph branches and support DUP

Upstream ggml-org#29622 adds a mixed token/embd branch to every input
embedding graph through ggml_build_forward_select(). Its nodes are
not flagged for compute, but the backend translated them anyway,
and the DUP in that branch was unsupported, so the scheduler split
the graph and passed the embeddings across the split with a fixed
token count. The first single-token decode then failed
(test-thread-safety on CPU and GPU).

Build the OV model from the compute nodes only, and translate a
same-type contiguous DUP like CONT so the graph stays on one backend.

* ggml-openvino: make inp_scale_rows token dim dynamic

ggml-org#29622 also moves the per-token embedding scale (gemma3, gemma3n,
gemma4) into a new [1, n_tokens] input. Give it a dynamic token dim
and pad it per chunk on the static (NPU) path.

* ggml-openvino: skip GPU MUL_MAT op tests with unbound Q4_1/Q4_K weights

Op tests build Q4_1/Q4_K weights as u4 with an f16 zero point. The GPU
plugin fails to compile that form for some row counts with "clFinish,
error code: -5 CL_OUT_OF_RESOURCES", which aborts test-backend-ops on
the MUL_MAT cases added in ggml-org#29869 (e.g. m=1000, n=2, k=1024). Model
weights use a u4 zero point and are not affected.

Report these cases as unsupported on GPU until the plugin is fixed.
Op tests check support before allocating, so the check matches unbound
weights only; model loading probes with a dummy buffer and keeps its
weights on the GPU.

* ggml-openvino: create FILL in the output type

translate_fill always built an f32 constant, so an f16 FILL produced
f32 data and the copy back overran the f16 output buffer. Use the
output type for the constant.

* ggml-openvino: reject CONCAT with a quantized type

Quantized inputs are dequantized when translated, so the backend cannot
write a quantized CONCAT output. Report it as unsupported, as for CPY
to a quantized type.

* ggml-openvino: handle the single recurrent state gather of build_rs

ggml-org#29856 changed build_rs to gather all recurrent states with one GET_ROWS
on the s_copy leaf and take the ubatch and extra states as views of it.
The stateful path matched only the previous form, a GET_ROWS per view of
s_copy, so Qwen3.5 failed with stateful execution on CPU and GPU
("is_axis_valid(axis, r)" in a Concat).

For a single-slot cache, treat the GET_ROWS on the s_copy leaf as the
active-state gather, keep the rank-4 layout of reshapes that read a view
of it, and map the copy of the empty extra-state view to the single-slot
remainder writeback. Do not warn about the dynamic dim of empty views.

* openvino: align eltwise operand ranks to work around a GPU-plugin defect

* openvino: match the MoE fusion on the rank-3 stateful graph

* ggml-openvino: do not unsqueeze an RMS norm output in AlignEltwiseOperandRanks

The pass unsqueezes the lower-rank operand of an Add/Multiply/Subtract
whose operand ranks differ. In gemma-3 the lower-rank operand of the
post-attention residual add is the norm output, and unsqueezing it makes
the GPU plugin compute the layer wrongly: gemma-3 returns empty answers
on GPU with stateful execution.

Skip the rewrite when the lower-rank operand is an RMS norm output.

* docs : update OpenVINO validated models

---------

Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com>
(cherry picked from commit b9a5a00)
@peteklapka

Copy link
Copy Markdown

Kudos. MTP finally does something on an M2 Ultra. Took Qwen3.8 27B Q8 from 21tok/s at ~20k ctx to 30tok/s. Loving the uplift.

thelittlefireman pushed a commit to thelittlefireman/llama.cpp that referenced this pull request Oct 12, 2026
* metal : few-row MMA mat-mul and batched copies for speculative decoding

Speculative decoding verifies a few draft tokens per step. Without the tensor API, Metal ran these mat-muls with the mat-vec kernels, whose time grows with every src1 row, so DFlash2 decoding on an M3 Ultra was slower than serial decoding.

- add mat-mul kernels for 2..16 src1 rows on 8x8 simdgroup matrices: each weight is dequantized once for all rows, and the simdgroups of a threadgroup split K. Q4_0, Q8_0 and Q5_K have their own kernels, F32, F16, Q4_1, Q5_0, Q5_1, Q4_K and Q6_K use a generic path over the 16-weight dequantizers, and Q4_0 at 2 rows uses a 2-row variant of the mat-vec kernel
- use them only on MTLGPUFamilyApple7+ without the tensor API, from the row count at which they beat the mat-vec kernels on an M3 Ultra (F32: 6, F16, Q4_K, Q5_0, Q5_1: 3, other types: 2)
- fusion table: MUL_MAT + ADD adds a same-shape residual in the MMA store, and up to 16 adjacent same-layout f32 copies between the same two tensors run as one dispatch
- the fusion checks and ggml_graph_optimize take the device props, so the reorder packs MUL_MAT + ADD only on devices that can fuse it, at every src1 row count
- views do not count toward GGML_METAL_FUSION_MAX when the reorder packs a group, so 16 recurrent state snapshot copies with views between them stay one group
- the encoder checks the inner nodes of a fused group for concurrency, tracks written views by their extent, and does not count the destination of a CPY as a read
- CONCAT splits long rows across threadgroups when there are few rows
- tests: few-row MUL_MAT, MUL_MAT_ADD, CPY_BATCH and CONCAT cases in test-backend-ops (with a prepare_graph hook for the copy order), test-metal-graph-optimize, test-metal-cpy-batch-alias

* metal : remove the CPY_BATCH fusion and the memory range changes

Remove the batched copy fusion with its kernel and tests, and revert the
memory range changes, as suggested in review. The memory ranges, the
graph reorder and the CPY encoder are again the same as on master.

* cont : clean-up

* cont : drop has_tensor gate

* cont : clean-up operand/residual logic

* cont : drop Q4_0 ne11=2 special-case

* cont : add kernels/mul_mv_mma.metal

* cont : consolidate mma pipeline selection logic

* cont : decouple fusion logic from device props

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
thelittlefireman pushed a commit to thelittlefireman/llama.cpp that referenced this pull request Oct 12, 2026
* ggml-openvino: skip unselected graph branches and support DUP

Upstream ggml-org#29622 adds a mixed token/embd branch to every input
embedding graph through ggml_build_forward_select(). Its nodes are
not flagged for compute, but the backend translated them anyway,
and the DUP in that branch was unsupported, so the scheduler split
the graph and passed the embeddings across the split with a fixed
token count. The first single-token decode then failed
(test-thread-safety on CPU and GPU).

Build the OV model from the compute nodes only, and translate a
same-type contiguous DUP like CONT so the graph stays on one backend.

* ggml-openvino: make inp_scale_rows token dim dynamic

ggml-org#29622 also moves the per-token embedding scale (gemma3, gemma3n,
gemma4) into a new [1, n_tokens] input. Give it a dynamic token dim
and pad it per chunk on the static (NPU) path.

* ggml-openvino: skip GPU MUL_MAT op tests with unbound Q4_1/Q4_K weights

Op tests build Q4_1/Q4_K weights as u4 with an f16 zero point. The GPU
plugin fails to compile that form for some row counts with "clFinish,
error code: -5 CL_OUT_OF_RESOURCES", which aborts test-backend-ops on
the MUL_MAT cases added in ggml-org#29869 (e.g. m=1000, n=2, k=1024). Model
weights use a u4 zero point and are not affected.

Report these cases as unsupported on GPU until the plugin is fixed.
Op tests check support before allocating, so the check matches unbound
weights only; model loading probes with a dummy buffer and keeps its
weights on the GPU.

* ggml-openvino: create FILL in the output type

translate_fill always built an f32 constant, so an f16 FILL produced
f32 data and the copy back overran the f16 output buffer. Use the
output type for the constant.

* ggml-openvino: reject CONCAT with a quantized type

Quantized inputs are dequantized when translated, so the backend cannot
write a quantized CONCAT output. Report it as unsupported, as for CPY
to a quantized type.

* ggml-openvino: handle the single recurrent state gather of build_rs

ggml-org#29856 changed build_rs to gather all recurrent states with one GET_ROWS
on the s_copy leaf and take the ubatch and extra states as views of it.
The stateful path matched only the previous form, a GET_ROWS per view of
s_copy, so Qwen3.5 failed with stateful execution on CPU and GPU
("is_axis_valid(axis, r)" in a Concat).

For a single-slot cache, treat the GET_ROWS on the s_copy leaf as the
active-state gather, keep the rank-4 layout of reshapes that read a view
of it, and map the copy of the empty extra-state view to the single-slot
remainder writeback. Do not warn about the dynamic dim of empty views.

* openvino: align eltwise operand ranks to work around a GPU-plugin defect

* openvino: match the MoE fusion on the rank-3 stateful graph

* ggml-openvino: do not unsqueeze an RMS norm output in AlignEltwiseOperandRanks

The pass unsqueezes the lower-rank operand of an Add/Multiply/Subtract
whose operand ranks differ. In gemma-3 the lower-rank operand of the
post-attention residual add is the norm output, and unsqueezing it makes
the GPU plugin compute the layer wrongly: gemma-3 returns empty answers
on GPU with stateful execution.

Skip the rewrite when the lower-rank operand is an RMS norm output.

* docs : update OpenVINO validated models

---------

Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com>
Achilles1089 added a commit to Achilles1089/llama.cpp that referenced this pull request Oct 15, 2026
…gml-org#29869 + ggml-org#30065, without their MUL_MAT+ADD fusion); stop accepting draft tokens at EOG (ggml-org#29638); graph capture flag fix (ggml-org#29390)

The 2..16-row mat-muls that MTP verify produces ran on mat-vec kernels whose time grows with every row. Measured on
the M1 Ultra with Raven's Qwen3.8-27B Q8_0 shapes: 8 rows now cost 1.6-1.9x one row (was 3.3-3.8x), matching
upstream within 1-3%. End to end (temp 1.0, 512 tokens): at --spec-draft-n-max 6 --spec-draft-n-min 1
--spec-draft-p-min 0.6, code echo 34.9 -> 54.7 tok/s and tool JSON 35.2 -> 54.9, prose level (25.6 -> 25.9).
Greedy output is identical to the previous build, with and without drafting. test-backend-ops MUL_MAT: 1163/1163.
The capture fix: the -1 capture counter was passed as a bool, so every node of every graph got debug-group markers.
gagallo7 pushed a commit to gagallo7/qvac-fabric-llm.cpp that referenced this pull request Oct 16, 2026
* metal : few-row MMA and batched copies for speculative decoding

- few-row simdgroup MMA in kernels/mul_mv.metal for 2..16 src1 rows: Q4_0, Q8_0 (up to 8 rows) and Q5_K kernels, and a generic path over the 16-weight dequantizers for F32, F16, Q4_1, Q5_0, Q5_1, Q8_0 (9..16 rows), Q4_K and Q6_K; a 2-row Q4_0 variant of the mat-vec kernel. The kernels fill simdgroup matrices per lane, so they run only on GPUs with native simdgroup matrices (MTLGPUFamilyApple7+) and without the tensor API
- two fusion table entries: MUL_MAT+ADD writes the residual sum from the MMA store, and up to 16 adjacent same-layout f32 copies between the same two tensors run as one dispatch (kernel_cpy_batch in kernels/quantize.metal). The table checks, ggml_metal_fusion_next/max and ggml_graph_optimize take the device properties, so graph_optimize packs MUL_MAT+ADD only on devices that fuse it
- graph_optimize packs what the table fuses, so it now packs runs of same-layout f32 copies and chains MUL_MAT+ADD with the table patterns after it; the original change packed only {MUL_MAT, ADD} and never packed copies. A pack may span more than 16 graph nodes when views sit between its nodes (views do not count toward GGML_METAL_FUSION_MAX), so a run of up to 16 conv state snapshot copies stays one batch; the copy batch entries are listed longest first, so the encoder checks no shorter batch after a match
- the graph reorder records what every node of a pack writes, not only the last node, so a reader of an inner output of a chained pack (the MUL_MAT+ADD sum before RMS_NORM+MUL, the copies into the first of two batches, the GATED_DELTA_NET output of a GATED_DELTA_NET+CPY pack) is not hoisted above it; this replaces the list of extra outputs of the topk-moe pack
- ggml_metal_mul_mat_use_mma excludes the FWHT (hadamard) mat-muls, which the encoder takes first, so the table packs MUL_MAT+ADD only for mat-muls the MMA store runs
- the fusions are skipped when the sum overlaps a mat-mul input or the residual (other than in place), when copy destinations overlap, or when a copy reads where another copy of the batch writes
- CONCAT, typed and block-quantized, splits long rows across threadgroups when there are few rows
- view-aware concurrency ranges in the encoder: written views are tracked by their extent and CPY does not count its destination as a read; the graph reorder keeps whole-tensor ranges, so graphs with the same nodes get the same order and allocation
- the radix TOP_K of the original change is upstream (2a74817) and is not duplicated
- M5, tensor API off, 5120-row Q4_0 (K 4096, 25600), Q5_K (K 6144) and Q8_0 (K 17408) weights, 2..16 rows: 0.36x-0.99x of the base kernel time, except Q8_0 at 2 rows (1.08x)
- tests: MUL_MAT_ADD, CPY_BATCH, positive-src1 MUL_MAT, MUL_MAT with K not a multiple of 64, CONCAT with few long rows, Metal copy-alias regression, Metal concurrency ranges, MUL_MAT+ADD packing per device and a reorder independent of written view sizes, graph reorder readers of chained packs and copy runs of 4..16 snapshots kept whole

* spec : DFlash2 draft vocabulary and copy-free dynamic conv

- add --spec-draft-vocab: a DFlash2 drafter computes its head over the given token id ranges only; the target still verifies over the full vocabulary
- pass the draft token ids as an I32 graph input, map the selector candidates with get_rows and scatter the logits with set_rows
- use head row views only when the head has no Hadamard rotation or LoRA delta and keeps its row layout: llama_model records the weight buffers of buffer types other than the device default or a host type (CPU_REPACK, AMX, KleidiAI, a row split, the CUTLASS repack), whose rows a view at a multiple of nb[1] does not address
- reject invalid ranges, ranges with fewer ids than the selector top-k, drafters that are not DFlash2 or have a reduced head, and other speculative types before any speculative state changes
- grow the DFlash2 graph node budget with the number of ranges
- feed the DFlash2 dynamic conv strided views instead of three CONTs per call
- tests: test-dflash-draft-vocab, test-arg-parser range syntax, REPEAT_MUL_VIEW, I32 GET_ROWS; test-dflash-loader covers the draft-vocab graph, a Q4_0 head (the ranges are ignored when the CPU repacks it), rejected, undersized and reduced-head ranges and one-token ranges over the whole vocabulary, and moves to the tests that use internal symbols; docs: speculative.md, cli and server README

* qwen35 : feed the GDN output projection a 2D input

The projection output then goes straight into the residual add, so the Metal MUL_MAT+ADD fusion applies to all 48 GDN layers.

* metal : pack MUL_MAT+ADD in the graph reorder for any number of src1 rows

The graph reorder packed a MUL_MAT with its residual ADD only when the
few-row MMA kernels take the mat-mul: 2 to 16 src1 rows, a K step that
depends on the row count for Q8_0, and not the 2-row Q4_0 kernel. A graph
reserved at n_ubatch rows and a decode graph of a few rows therefore got
different packs and a different node order. They have the same nodes, so
ggml_gallocr, which keeps the allocation per node position, had to
reallocate. On a GPU without the tensor API,
test-recurrent-state-rollback-dsv4 and -kimi-k3 aborted on an unexpected
graph reallocation with GGML_SCHED_NO_REALLOC; without it, the scheduler
re-reserves the graph each time the order changes.

The structural check now reads only the device, the Hadamard hint, the
types and whether the added tensor is a weight, never shapes or strides,
so a mat-mul and an add of its shape get the same order at every batch
size. A bias lives in a weights buffer, so its add stays free to run
next to independent nodes, such as the Q, K and V projections, as
before. The full check of the encoder starts from the structural one, so
the encoder fuses only pairs that the reorder packs, and still only
few-row mat-muls with a same-shape residual. A mat-mul with a Hadamard
hint is neither packed nor fused.

test-metal-graph-optimize checks that the reorder of a MUL_MAT with a
residual, a one-row bias weight or a view of one is the same at 1, 4, 9,
64 and 512 src1 rows, that the residual add is packed and the bias adds
are not, and that the encoder check fuses a same-shape residual but not
one that is a weight. The residual case fails with the old rule, the bias
and encoder cases without the weights check, the view case without its
lookup of the viewed tensor.

(cherry picked from commit 7371f38)
(cherry picked from commit 0bab9c7)
(cherry picked from commit f25ec33)

b11514 rebase: Split pick: the Metal half is upstream a3a1c47
(ggml-org#29869) as on HEAD; the duplicate MMA kernels, CPY_BATCH and its tests
are dropped. The DFlash2 draft-vocab spec half and the REPEAT_MUL_VIEW
test are kept (see DECISIONS-b11514.md).

Assisted-by: Claude Opus 5.5
Signed-off-by: Guilherme Gallo <guilherme.gallo@collabora.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Apple Metal https://en.wikipedia.org/wiki/Metal_(API) ggml changes relating to the ggml tensor library for machine learning highlight Changes that will be highlighted in the next release notes testing Everything test related

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants