metal : add fa-vec tunings for M1 Ultra (#28088)

• metal : add fa-vec tunings for M1 Ultra

• metal : move M1 Ultra tunings after M1 Max section

• metal : remove duplicate blank line

CUDA: XOR swizzle flash attn K,V smem fp16 tiles (#25635)

• CUDA: XOR swizzle flash attn K,V smem fp16 tiles

Signed-off-by: ynankani ynankani@nvidia.com

• Fix use 64bit generic pointer instead of 32bit shared pointer

Signed-off-by: ynankani ynankani@nvidia.com

• fix shared memory race in FA on DGX Spark

• Handle corener case

Signed-off-by: ynankani ynankani@nvidia.com

• Add swizzle test cases and gate sync for swizzled path only

Signed-off-by: ynankani ynankani@nvidia.com

• gate CUDA PTX

Signed-off-by: ynankani ynankani@nvidia.com

• offset calculation specific for swizzle branch

Signed-off-by: ynankani ynankani@nvidia.com

• Reafctor code

Signed-off-by: ynankani ynankani@nvidia.com

• Refactor FA swizzle ldmatrix if/else into helpers (K row/col, V offset)

Signed-off-by: ynankani ynankani@nvidia.com

• rebase and update test case args

Signed-off-by: ynankani ynankani@nvidia.com

• Allow swizzle for non-pow2 shapes, for which nbatch_2%32==0

Signed-off-by: ynankani ynankani@nvidia.com

Signed-off-by: ynankani ynankani@nvidia.com

metal : add concat support for quantized types (#28116)

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731

AVX2: Speed up large batch size prompt processing of IQ models (#27402)

• Batched gemm for grid IQ quants

Style updates and a bit more performance

Clean up comments

Move code around

Vectorize IQ panel decode, lower threshold for speedup

IQ panel: single-source gather layout, gate bias, vectorize interleave

Add ggml_gemm_iqp_8x8_q8_K_p4 kernel, remove gather buffer

Move IQ panel code out of repack into iqp.cpp, clean up comments

Another comment sweep

• Add myself as iqp.* codeownder

• Remove ggml_cpu_iqp_scratch_offset and ggml_cpu_iqp_src1_conv_size

• Renaming and moving

• The other half of renaming and moving

• Move macros and ggml_cpu_iqp_mul_mat_id_min_batch definition

• Update ggml/src/ggml-cpu/iqp.h

Co-authored-by: Georgi Gerganov ggerganov@gmail.com

• Add iqp_rows work buffer

• Revert "Add iqp_rows work buffer"

This reverts commit 4255429.

• Add NUMA fallback

• Add 10 row batch tests for IQP coverage on all grid IQ types

• Swap assert for return false in support check

• Move IQP mul_mat_id test

Co-authored-by: Georgi Gerganov ggerganov@gmail.com

kv-cache : optimize restoring non-contiguous cells (#27991)

• kv cache : batch state restore scatter reads per contiguous run

When restoring state into non-contiguous destination cells (e.g. a
prompt-cache snapshot into a fragmented ring), state_read_data issued
one small copy per KV cell - ~1.4M copies of a few KiB each for a
40k+ token restore, taking 25-63 s on the CUDA backend.

The snapshot stores cell rows in cell order, so a maximal run of
consecutive destination indices maps to one contiguous block and can
be restored with a single copy. Precompute the runs once and use them
in all three scatter loops (K, V, transposed V). Byte-identical.

The on-device reader copies with a byte cursor when the read and
write chunking differs, so the batched reads are safe for it as well.
Batching makes equal tensor counts with a different split reachable
(save ranges [2,1] vs restore runs [1,2]); the next commit teaches the
reader's 1:1 path to fall back to the byte cursor in that case.

Verified in a production setup: 1,363,616 copies / 25-63 s -> 224
copies / 221-424 ms for the same restores (42,603 cells, 4 runs).

Assisted-by: Claude Code (unsloth/qwen3.8-27b)

• context : fall back to the byte cursor when read and write chunking differ

the on-device reader copies saved state back with a 1:1 copy by tensor
index whenever the write and read sides recorded the same number of
tensors, guarded by a per-tensor size assert.

equal tensor counts do not imply equal chunking: a state restore may
batch its reads per contiguous run of destination cells while the save
used per-range reads, so both sides can record two tensors that split
the same data differently, and the assert aborts in all builds.

compare the per-tensor sizes and only take the 1:1 path when the
chunking actually matches, otherwise fall through to the existing
byte-cursor copy. both sides enumerate the same logical data in the
same order, so the cursor copy is well-defined across tensor
boundaries.

Assisted-by: Claude Code (unsloth/qwen3.8-27b)

• tests : cover state restore scatter reads on host and on-device paths

decode the same prefix on two sequences, interleaving the seq 0 cells
between the seq 1 cells, so the seq 1 cells are isolated from each
other in the kv cache (three cells, two saved ranges). save the seq 1
state, free the interleaved seq 0 cells, and restore: the destination
is then non-contiguous (two runs), and the restore-side chunking has
the same tensor count as the save-side with a different split, so the
scatter path is batched per contiguous run and the on-device reader's
byte-cursor fallback is exercised.

the restored state is saved again on the host and compared byte for
byte with the first save: the blob is serialized in sequence cell
order, so the two saves are identical if and only if the scatter
restore wrote exactly the same KV content. this documents the
byte-identical guarantee of the run-batched scatter reads.

one test per io backend: the host (CPU) path and the on-device path.

Assisted-by: Claude Code (unsloth/qwen3.8-27b)

opencl: tune the quant paths for Intel Xe-LP GPUs to improve its TG and PP performance (#26438)

• opencl: Q4_K/Q5_K mul_mv N_DST 4->8 on Intel for 2x activation reuse

• opencl: Q4_K mul_mm 8x8 tile fot Intel

• opencl: Q5_K mul_mm 8x8 tile for Intel

• opencl: Q4_K mul_mv N_DST 8->16 for Intel

webgpu : avoid crash when offset is not multiple of 4 in WebGPU ggml_backend_tensor_get() implementation (#28045)

• webgpu : avoid crash when offset is not multiple of 4 in WebGPU ggml_backend_tensor_get() implementation

• chore : improve code readability

Co-authored-by: Sigbjørn Skjæret sigbjorn.skjaeret@huggingface.co

Co-authored-by: Stanisław Szymczyk sszymczy@gmail.com
Co-authored-by: Sigbjørn Skjæret sigbjorn.skjaeret@huggingface.co

ROCm: add radix TOP_K for long rows (#27466)

• ROCm: add radix TOP_K for long rows

metal : add fa-vec tunings for M1 (#28078)

CUDA: extend MOE fusion to specdec, earlier MOE glu fusion and topk-router fusion were restricted to 1 token (#27621)

• CUDA: extend MOE fusion to specdec, earlier MOE glu fusion and topk-router fusion were resticted to 1 token

Signed-off-by: ynankani ynankani@nvidia.com

• Address review comments

Signed-off-by: ynankani ynankani@nvidia.com

• Add SWIGLU_CLAMP case to multi-token moe fusion

Signed-off-by: ynankani ynankani@nvidia.com

Signed-off-by: ynankani ynankani@nvidia.com