* metal : fix half-idle simdgroup in kernel_mul_mv_iq3_xxs_f32 for ne00 < 1024
* metal : keep N_R0_IQ3_XXS = 4, dispatch a separate 8-row split kernel for ne00/32 < 32
The plain kernel is unchanged from master (4 rows per simdgroup, one thread per
chunk). The row-split mapping now lives in a separate kernel_mul_mv_iq3_xxs_f32_split
instantiation with N_R0_IQ3_XXS_SPLIT = 8, and the host selects it only when
ne00/32 < 32 and divides 32, so wide matrices keep the master kernel bit for bit.
* metal : select the iq3_xxs row split with a function constant instead of a separate kernel
* metal : support n_kv_max sparse mask hint in flash attention vec kernel
- add kernel_flash_attn_ext_vec_idx: compacts finite mask entries into
a per-row index list (Hillis-Steele scan, one threadgroup per row)
- extend vec FA kernel with optional sparse index gathering (FC slot 5)
- add host-side gate: sparse path when n_kv_max > 0, mask present,
supported head sizes / KV types, n_kv_max <= 4096
- new buffer region extra_idx for the index list
- pipeline getter extended with has_sparse param
- add test cases: head sizes, quant types, nb>1, nr23 variants,
sinks, ALiBi, softcap, permute, v_view_of_k, no-mask fallback
Note: multi-row (nb*nr23[1] > 1) cases still failing - rid mapping
in the store phase needs revisiting for the sparse path.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* metal : fix sparse flash attention row addressing
- kernel_flash_attn_ext_vec_idx: mask param is half* but nb31 is a byte
stride, so the per-row mask offset was scaled by 2x; cast to char*
before applying the byte strides
- kernel_flash_attn_ext_vec: sparse pidx param is char* so the per-row
element offset was under-scaled by sizeof(int); scale it by sizeof(int)
to get the correct byte offset
- fixes the multi-row (nb*nr23[1] > 1) sparse flash attention failures
Assisted-by: pi:llama.cpp/DeepSeek-v4-0731
* cont : use sparse vec FA for prefill
* metal : single-pass flash attention sparse index compaction
The idx kernel previously read the mask row twice: once to count the finite
entries (for the prefix scan) and again to recover their positions. Since the
kernel is memory-bound, this doubled the mask traffic.
Keep the finite positions in a per-thread register array during the count
pass and write them out directly, avoiding the second mask read. A dense
mask with more than NLOCAL finite entries in a slice falls back to re-reading
the mask to write the remaining positions.
Assisted-by: pi:llama.cpp/DeepSeek-v4-0731
* tests : add perf cases for sparse flash attention prefill
Measure the sparse vec FA kernel across KV sizes, n_kv_max hints and batch
sizes. Run with:
./build/bin/test-backend-ops -b MTL0 -o FLASH_ATTN_EXT -p "n_kv_max=[1-9]" perf
Assisted-by: pi:llama.cpp/DeepSeek-v4-0731
* qwen4 : enable sparse attention
* cont : adjust nsg
* cont : sync test-backend-ops
* cont : disable Qwen4 for now
* cont : clean-up + tests
* metal: WIP chunked SSD SSM_SCAN kernels for multi-token prefill
* metal: drop scalar SSD path; MMA + sequential tail
* drop WIP ssm scan test noise
* remove state_from_dst and rename CS and NSG constants
* remove unrelated added whitespace padding
* added clarity to mma_tokens calculation
* added clarity to use_mma bool checks
* added comments to metal ssd op constants for clarity
* reserve K tokens for sequential kernel rollback snapshots
* reset concurrency between mma and seq tail
* remove print args no longer used
* fixed comment to no longer point to specific line
* add FC_SSM_SCAN so seq path skips token offlset unless it's mma tail
* added changes to new ssm.metal for rebase after ggml-metal.metal refactor
* specialize ssm_scan tail with a template instead of a function constant
---------
Co-authored-by: dpantaleoni <dominikpantaleoni@gmail.com>
Co-authored-by: forforever73 <690105611@qq.com>
* metal : per-device tuned (Q, NE) for flash-attn vec (#25750)
* rebase Q-generic FA vec body from 01dc93607 (#23114)
* add 53 f16 (Q,NE) flash-attn vec instantiations (vec 80 -> 133)
* add FA vec (Q,NE) tuning table + dispatch wiring + SMEM cap fallback
* add FA vec (Q,NE) perf sweep
* fill tuning result
* fold family table into a per-family representative SKU
* refactor tuning result format
* extend FA vec tuning to quantized KV caches
* sync fa vec tuner bucketing with runtime, use pointwise tuning regret
* update tuned table
* format and cleanup
* prefix fa_vec tuning procs with ggml_backend_metal_tuning_, drop unused fa_vec_override_active
* add device id -> token lookup for the offline tuning tool
* add ggml-metal-tuning skeleton
* add op-agnostic perf cell + median timing for the tuner
* add FA-vec graph build + tensor init to the tuner
* tools : add FA-vec (Q,NE) sweep, compression and table emit
* cool down and re-measure the dirty window on thermal drift
* test-backend-ops : replace the FA vec tune mode with a bounded (Q,NE) slice
* tools : document the Metal tuner, point the table comment at it
* abort on unknown KV type, single-source fa_vec_legal_ne
* cleanup
* honor -o in the FA vec (Q,NE) slice
* retune FA-vec (Q, NE) under a pointwise no-harm gate
* cont : add fa-vec tunings for M1 Pro, M2 Ultra, M5 Max
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* metal: dequantize q8_0 KV to f16 before flash attention
Add a preprocessing pass for GGML_OP_FLASH_ATTN_EXT on the Metal backend:
when the KV cache is quantized (Q8_0 for now), dequantize K and V into a
contiguous F16 scratch buffer and run the existing F16 flash attention
kernels on it, instead of the in-kernel dequantization path.
- new kernel kernel_flash_attn_ext_dequant_to_f16<block_t, QK, deq_t4x4>:
one thread per quant block (K then V), stride-aware so permuted KV is
supported; instantiated for Q8_0 (extending to Q4_0/Q4_1/Q5_0/Q5_1 is
one instantiation + one gate case)
- the gate is type-only: dequantize whenever the KV is quantized,
regardless of head sizes, GQA ratio or n_kv; the attention kernels
themselves are untouched
- the F16 copies live in the op's own scratch allocation
(ggml_metal_op_flash_attn_ext_extra_dequant_f16); the KV pad kernel
reads the dequantized buffers when the path is active
- the FA pipeline getters gain a use_f16_kv flag selecting the existing
f16 kernels and contiguous strides
- ref: https://github.com/ggml-org/llama.cpp/pull/25556
Verification (M2 Ultra):
- test-backend-ops test -o FLASH_ATTN_EXT: 4798/4798 pass, including the
new q8_0 eval cases (decode/prompt, permuted, sinks+ALiBi+softcap,
kv=113 pad path, kv=16384)
- llama-perplexity on Qwen2.5-0.5B with -ctk q8_0 -ctv q8_0 matches the
f16 KV reference (PPL 1.0008 vs 1.0008)
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* metal : launch the FA KV dequant kernel separately for K and V
Simplify kernel_flash_attn_ext_dequant_to_f16: it now dequantizes a single
tensor (its own ne/nb and dst) with no is_v branching, and the op dispatches
it twice with the same pipeline - once for K and once for V. The kargs
struct shrinks to a single ne/nb set plus nblocks.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* metal : dequantize q4_0, q4_1, q5_0 and q5_1 KV to f16 before flash attention
The dequant pass now covers all quantized KV types supported by the Metal
flash attention kernels. The dequant kernel, kargs, scratch allocation and
dispatch are type-generic, so each type is one kernel instantiation plus one
gate case.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* metal : skip the redundant V dequant when V is a view of K
In MLA-based models, the V of the FA op is a view of K (the first ne20
elements of each K row); the dequantized V is then a view of the dequantized
K, so skip the second dequant dispatch, do not reserve the V scratch region,
and let the pad and attention kernels read V from the K F16 buffer with K's
strides. The detection follows the CUDA backend:
V->view_src && (V->view_src == K || (V->view_src == K->view_src && V->view_offs == K->view_offs))
Also fix the FA pipeline getters: ns10/ns20 are function constants baked into
the kernels and must be the actual K/V row widths as seen by the kernel. The
dispatch now passes them explicitly (nb11_attn/nb10_attn, nb21_attn/nb20_attn)
instead of the getters assuming contiguous F16 KV (ns20 = dv), which was wrong
when V is read from K with K's row pitch (e.g. 576 vs 512).
New test cases: 576/512 q8_0 (MLA shape, V is a view of K) at kv=113 (KV pad),
nb=1 (vec) and nb=64 (non-vec).
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* test : remove backend-specific wording from test-backend-ops comments
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* pi : avoid backend mentions in test-backend-ops comments
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* metal : rename the FA dequant_f16 identifiers to kv_f16
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* cont : clean-up
* cont : remove TODO
* metal: add TQ2_0 support
Add support for the GGML_TYPE_TQ2_0 (ternary, 2 bits per element) type in
the Metal backend.
Assisted-by: llama.cpp:DeepSeek-v4-Flash-0731
* cont : optimize mul_mv kernel
- float ops over integer ops
- precalculate sums
- hoist coef out of the inner loop
- contiguous y loads
llama.cpp:DeepSeek-v4-Flash-0731
* feat(silu_back): implemented silu_back op for f32
* fix(silu_back): removed redundant asserts in ggml-metal-ops.cpp function ggml_metal_op_silu_back.
- Implement GGML_OP_DSV4_HC_COMB, GGML_OP_DSV4_HC_PRE, and
GGML_OP_DSV4_HC_POST with SIMDgroup register and shuffle optimized kernels.
- Add Metal dispatch and support plumbing and test the production Sinkhorn
iteration count and embedding width.
Assisted-by: Codex
Co-authored-by: Thiago Padilha <thiago@padilha.cc>
* metal: fuse snake activation (mul, sin, sqr, mul, add)
Mirror the CUDA, Vulkan and CPU snake fusion: same matcher on the naive
5-op chain, same F32 contract on a and inv_b, same F32/F16/BF16 kernel
with F32 compute. Follows the Metal backend idioms: bf16 instantiation
gated behind GGML_METAL_HAS_BF16 and concurrency ranges checked on the
remaining chain nodes before encoding, as done by the bin fusion.
Covered by the existing backend-agnostic SNAKE_FUSE tests.
* metal: absorb snake fusion into ggml_metal_op_bin
Extract the matcher to ggml_metal_op_can_fuse_snake, mirroring the
Vulkan naming, and dispatch the fused path from ggml_metal_op_bin.
The encode loop switch is back to a single call per case.
Address review from ggerganov
* metal: fix indentation in ggml_metal_op_can_fuse_snake
* metal : add CONV_2D_DW (depthwise 2D convolution) support
* test : add perf cases for CONV_2D_DW
* metal : use 3D dispatch for CONV_2D_DW kernel
* metal : add channel-tiled CONV_2D_DW kernel for non-contiguous layouts
* metal : simplify CONV_2D_DW dispatch and trim comments
* metal : merge duplicate CONV_2D_DW pipeline getters
* tests : add F16 CONV2D_DW tests
* cpu : fix F16 kernel support for CONV_2D_DW
* tests : remove commented-out CONV_2D_DW test block
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* metal: add col2im_1d op (f32/f16/bf16)
Gather kernel mirroring the CPU/CUDA path: each output (t_out, oc)
reads its ceil(K/s0) source columns with an F32 accumulator, a single
write and no atomics. One thread per output element, 256 per
threadgroup.
* metal: check dst contiguity and type match in supports_op for COL2IM_1D
Align the GGML_OP_COL2IM_1D predicate with the CPU, CUDA, and Vulkan
backends: the kernel writes dst with linear indexing and assumes the
same type as src0, so supports_op must also require a contiguous dst
and op->type == op->src[0]->type.
* Update ggml/src/ggml-metal/ggml-metal.metal
Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com>
---------
Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com>
Reuse existing rope kernels with a function constant to toggle forward/backward
rotation, avoiding duplicate kernel code.
Assisted-by: pi:llama.cpp/Qwen3.6-27B
* metal : add f16 and bf16 support for concat operator
Extend the Metal backend concat operator to support f16 and bf16 tensor
types in addition to the existing f32 and i32 support.
- Template kernel_concat on type T with specializations for float, half,
bfloat, and int
- Add type-specific pipeline getter ggml_metal_library_get_pipeline_concat()
- Update device support check to allow f16 unconditionally and bf16 when
device supports bfloat16
- Update dispatch to select the correct kernel specialization by type
Assisted-by: pi:llama.cpp/Qwen3.6-27B
* metal : extend concat operator to support f16, bf16, i8, i16 and i64
Assisted-by: pi:llama.cpp/Qwen3.6-27B
* Make ggml_gated_delta_net take only the initial recurrent state (D, 1, n_seqs) and passes the snapshot count K as an op parameter instead of inferring it from state->ne[1].
Remove the padding hack and copy all emitted snapshots into the recurrent cache with a single strided ggml_cpy
* Make GDN changes in all backends. Address review comments.
* Fix CI build errors
* spec: support MTP
* fix batch size
* rename files
* cont : simplify (#7)
* MTP: clean-up (#9)
* MTP: clean-up
* review: use llama_context_type instead of llama_graph_type
* review: remove llama_model_has_mtp
* review: fix convert issues
* convert: fix pycheck
* review: formatting
* use `mtp-` for identifying mtp models
* convert: fix mtp conversion
* mtp -> draft-mtp
* remove unused llama_arch
* add need_embd in speculative
* llama: allow partial seq_rm for GDN models for speculative decoding
Currently speculative checkpoint needs to restart from a checkpoint
after some draft tokens are not accepted, this leads to some wastage in
running the target again. This PR adds the ability to rollback upto
`draft_max` by storing the GDN intermediates.
* fix pending state
* vulkan: add GDN partial rollback
* meta: extend check to axis 1
* metal: add GDN partial rollback
Extend the gated delta net kernel to store intermediate states for
partial rollback support on the Metal backend.
- Add K (snapshot slot count) as a function constant
- Read input state from slot 0 of the 3D state tensor
- Write intermediate states to different slots during token loop
- For K=1, maintain backward-compatible single-slot behavior
Ref: https://github.com/ggml-org/llama.cpp/commit/8c05923630110223669f069af2000e9cf10c02bc
Assisted-by: llama.cpp:local pi
* delta_net_base: use ggml_pad instead of new_tensor
* review: add need_rs_seq
* review: rename part_bounded to n_rs
* review: deslop comments
* review: rename, add asserts
* server : adjust checkpoint logic (#11)
* server : adjust checkpoint logic
* cont : rm asserts
* server-context: fix early exit
* spec : fix compatibility with n-gram and add TODOs (#13)
* metal : cleanup
* llama : fix faulty bitwise check in recurrent memory
* server : disable RS-based MTP in combination with other spec types
* spec : add TODOs
* cont : fix comment
* cont : update comment
* common : fix logic for ngram + mtp compat
* llama-memory: enable checkpointing with partial rollback
* cont: add test-case for loading into a dirty ctx
* llama-memory-recurrent: clear rs_idx in clear
* download: fix mtp path
* llama-arch: fix enorm op
* docs: update docs
* conversion: fix type annotations
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* Optimize Metal Tensor API usage for matmul2d
Separates the Metal Tensor API (matmul2d) path in kernel_mul_mm into its own standalone kernel, gated by GGML_METAL_HAS_TENSOR.
The legacy simdgroup_matrix kernel is preserved under #else.
Previously both paths were interleaved via #ifdef blocks within a single kernel, forcing the tensor path to share the legacy kernel's data layout and threadgroup memory scheme. Splitting the kernel enabled memory and dispatch optimizations that weren't possible when the two paths shared code structure.
* cont : cleanup
* cont : cleanup
* cont : cleanup
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* nix: support unified apple-sdk
* Impl roll op for Metal
* Revert "nix: support unified apple-sdk"
This reverts commit abfa473360471532c547de8b202c780507924d4b.
* update ops.md
* update op docs
* initial Q1_0 Metal backend
* tuning q1_0 metal kernels
* add Q1_0 to test-backend-ops
* add Q1_0<->F32 copy test
* Apply suggestions from code review
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* Apply suggestions from code review
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* metal:add conv_3d backend
Rebased with master and resolved conflicts.
* Resolved issues related to changes in variable names
* kernel void kernel_upscale_bilinear_f32 was missing in my branch, added back, should pass all tests now
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* llama : enable chunked fused GDN path
* models : avoid Q and K repeats when using fused GDA
* cont : fix comment
Co-authored-by: Aman Gupta <amangupta052@gmail.com>
* cont : fix the fix
Co-authored-by: Aman Gupta <amangupta052@gmail.com>
* cont : fix
* metal : add GDN kernel (#20361)
* metal : add Metal backend for GGML_OP_GATED_DELTA_NET
Add a fused Metal kernel for the gated delta net recurrence op
(#19504), enabling GPU-accelerated inference for DeltaNet-based
models (Qwen3.5, etc.) on Apple Silicon.
Supports both GDA (scalar gate) and KDA (per-row gate) modes
with head_size 64 and 128. Unsupported configurations (head_size
32, non-contiguous tensors) gracefully fall back to CPU.
Performance: Qwen3.5-0.8B Q4_K_M on M4 Max
tg128: 170 -> 213 t/s (+25%)
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* metal : validate contiguity of all input tensors in supports_op
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* metal : add algorithm equivalence comment for GDA decay path
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* cont : unslop + optimize
* cont : clean-up
---------
Co-authored-by: Paul Flynn <paul@arkavo.com>
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
* CUDA: AR gated delta net improvements (#20391)
* Add FastDiv to gated_delta_net_cuda
* Shard columns across warps
This reduces register pressure (avoids spill for S_v = 128) and gives
the warp-scheduler more CTAs to schedule (thus hiding data-access
latencies).
* Remove unneded include in gated_delta_net.cu
* Improve comments
* Apply code-formating
* Make sharding HIP-compatible
1. Use ggml_cuda_get_physical_warp_size() to determine warp size flexibly
2. Add test with partial warp to test sum reduction on CUDA
* Remove fastdiv_s64, as we can treat neqk1 and rq3 as uint32_t
* Rename variables
* Enable GDN also for prefill, move TODO for chunked_GDN
* Actually remove the TODO from 206890897546bd16602c3b79394fd5ea09ef199f
* Get warp size at runtime
warp_size is not known at compile time in hip host code.
* Don't expose ggml_cuda_get_physical_warp_size on host
---------
Co-authored-by: uvos <devnull@uvos.xyz>
* llama : refactor llm_build_delta_net_base API
---------
Co-authored-by: Aman Gupta <amangupta052@gmail.com>
Co-authored-by: Paul Flynn <paul@arkavo.com>
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
Co-authored-by: Oliver Simons <osimons@nvidia.com>
Co-authored-by: uvos <devnull@uvos.xyz>
* add count equal for metal
* remove trailing whitespace
* updated doc ops table
* changed shmem to i32
* added multi tg and templating
* removed BLAS support from Metal docs
* Apply suggestions from code review
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* add memset to set dst to 0
* metal : cleanup
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* feat: Add a batched version of ssm_conv
This was done using Claude Code. It found a number of optimizations around
how the threads were organized, resulting in a huge performance boost!
Branch: Mamba2SSD
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* feat: Optimized SSM_SCAN kernel for metal
This used Claude Code and resulted in a modest performance improvement
while maintaining correctness.
Branch: Mamba2SSD
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* test: Add test-backend-ops perf tests for SSM_CONV
Branch: SSMKernelImprovements
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* test: Real representitive tests for SSM_CONV
Branch: SSMKernelImprovements
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* refactor: Use function constant for ssm_conv batch size
Branch: SSMKernelImprovements
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* test: backend op tests for ssm_scan from granite4 1b-h
Branch: SSMKernelImprovements
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* style: remove commented out templates
Branch: SSMKernelImprovements
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* feat: float4 version of ssm_conv_batched
Branch: SSMKernelImprovements
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* fix: Add missing ggml_metal_cv_free
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
---------
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* feat(wip): Port initial TRI impl from pervious work
The kernel does not work and is not optimized, but the
code compiles and runs, so this will be the starting point
now that the core op has been merged.
Branch: ggml-cumsum-tri
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* fix: Remove argument for constant val override
This was added in the original draft, but later removed. With this, the
kernel now passes tests.
Branch: ggml-cumsum-tri
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* feat: Move the ttype conditional to templating to avoid conditional in kernel
Branch: ggml-cumsum-tri
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* fix: Type fixes
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* feat: Add softplus for metal
Branch: ggml-cumsum-tri
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* feat: Add EXPM1 for metal
Branch: ggml-cumsum-tri
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* feat: Add FILL for metal
Branch: ggml-cumsum-tri
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* refactor: Branchless version of tri using _ggml_vec_tri_cmp as a mask
Branch: ggml-cumsum-tri
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* fix: Remove unused arguments
Branch: ggml-cumsum-tri
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
* refactor: Use select instead of branch for softplus non-vec
Branch: ggml-cumsum-tri
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
---------
Signed-off-by: Gabe Goodhart <ghart@us.ibm.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Foad Abo Dahood