opencl: add q5_0/q5_1 gemm and gemv kernels for Adreno (llama/24319)

* opencl: add q5_0 adreno support

* opencl: add q5_1 adreno support

* opencl: cosmetic fix

---------

Co-authored-by: Li He <redacted>

vulkan: ifdef eMesaHoneykrisp (build fix) (llama/24479)

Fixes build/CI after #24306.

ggml : bump version to 0.15.0 (ggml/1539)

vulkan: add fast path for contiguous buffer transfers (llama/23973)

vulkan: use medium matmul tile on Asahi Linux (llama/24306)

* vulkan: use medium matmul tile on Asahi Linux

* vulkan: switch Apple detection to Honeykrisp driver id

Remove padding and multiple D2D copies for MTP (llama/24086)

* 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

CUDA: Fix ssm_scan_f32 data-races (llama/24360)

* Add missing syncthreads before resuing cub_temp_storage

__syncthreads() is required before being allowed to resue TempStorage
smem:
https://nvidia.github.io/cccl/unstable/cub/api/classcub_1_1BlockLoad.html#_CPPv4I0EN3cub9BlockLoad4LoadEv20RandomAccessIteratorRA14ItemsPerThread_1Ti

* Add one more missing __syncthreads

Could also double-buffer, but alternative is to simply ensure all
threads have read smem* before writing to it again in the next loop
iteration

* Remove unused smem from ssm_scan_f32

vulkan: reduce iq1 shared memory usage for mul_mm (llama/24287)

vulkan: add `v_dot2_f32_f16` support in matrix-matrix multiplication and Flash Attention (llama/24123)

* vulkan: add support for valve fp16 dot2 extension

* use macro for dot2 path choice

* properly check for the feature

* add dot_product abstraction to reduce preprocessor branching

ggml : add GGML_OP_COL2IM_1D (llama/24206)

* cpu: add GGML_OP_COL2IM_1D

Add the overlap-add (scatter-add) step of a 1D transposed convolution.
A ConvTranspose1d factorizes as a GEMM followed by col2im: a weight
pre-permuted to [IC, K*OC] is contracted against the [IC, T_in] input
with mul_mat to produce a column matrix [K*OC, T_in], and col2im_1d
scatters those columns back into the [T_out, OC] signal, with
T_out = (T_in - 1)*s0 + K - 2*p0.

Keeping the contraction as a plain mul_mat leaves the heavy work on the
optimized (and quantizable) matmul kernels, so col2im_1d only does the
cheap overlap-add.

CPU uses a gather formulation parallelized over output channels,
supporting F32, F16 and BF16 with an F32 accumulator.

* tests: add backend coverage for GGML_OP_COL2IM_1D

Add test_col2im_1d next to the conv_transpose_1d cases, covering F32,
F16 and BF16 across eight geometries: the canonical kernel = 2*stride
DAC upsampling shape, overlap, no overlap, cropping (p0 = 1 and
p0 = stride/2), kernel < stride with zeroed gaps, kernel not a
multiple of stride, and a single column unfold.

Perf mode gets three real vocoder stage shapes reporting memory
bandwidth. max_nmse_err relaxes to 5e-4 for F16 and BF16.

* cpu: harden GGML_OP_COL2IM_1D

ggml_col2im_1d validates s0, oc, p0 and input contiguity at graph
build time, before the oc division, protecting every backend at once.
The kernel asserts the contiguity its flat indexing assumes and its
doc states the full output length including the crop term.

The kernel parallelizes over the time axis: the split stays balanced
down to OC = 1, where the previous channel split was single threaded.
Values are bit identical on the three real vocoder chains, two out of
three improve.

* tests: extend the GGML_OP_COL2IM_1D grid

The eval grid grows to eleven geometries: OC = 1 (mono output stage),
K = 1 with stride > 1 (sparse scatter, every gap position zeroed) and
a crop down to T_out = 2 where all the gather bounds act at once.

* tests: add col2im_1d equivalence test

tests/test-col2im-1d.cpp proves mul_mat + col2im_1d matches the
native ggml_conv_transpose_1d on the CPU backend, F32 bit exact, F16
and BF16 through casts of the column matrix. test-backend-ops cannot
cover this for a CPU only op since the CPU backend is its own
reference there.

* rpc: bump protocol patch version for GGML_OP_COL2IM_1D

GGML_OP_COUNT goes from 96 to 97 with the new op, which trips the
static_assert in ggml-rpc.h. Bump RPC_PROTO_PATCH_VERSION since the
op is appended and no existing op code shifts.

ggml-cpu : fix rms_norm_back wrong output under in-place aliasing (llama/24305)

* ggml-cpu : fix rms_norm_back wrong output under in-place aliasing

* cont : clean-up comment

---------

Co-authored-by: Georgi Gerganov <redacted>

Remove case for GGML_TYPE_Q4_K in mvvq.cu (llama/23528)

ggml-webgpu: Add clang-format job (llama/24308)

* Add clang-format job

* try local formatting

ggml-webgpu: Improve prefill speeds for k-quants + refactor matmul for Q4/Q5/Q8 and k-quants (llama/24225)

* ggml-webgpu: Improve prefill speeds + refactor matmul for quants

* Fixes for editroconfig checker

Handle buffer overlap / buffer aliasing for concat operator (llama/24000)

* Only run webgpu CI on my fork

* Add webgpu only workflow

* handle buffer overlap case for concat operator

* restore build-webgpu.yml

Co-Authored-By: Claude Sonnet 4.6 <redacted>
* Run clang-format

* Update ggml/src/ggml-webgpu/wgsl-shaders/concat.wgsl

---------

Co-authored-by: Claude Sonnet 4.6 <redacted>
Co-authored-by: Reese Levine <redacted>

Implement 2D workgroups for scale, binary, and unary ops (llama/24044)

* Only run webgpu CI on my fork

* Add webgpu only workflow

* Implement 2d workgroups for more operations

* fix

* Fix type

* Move back to global_invocation_id

vulkan: Use cm2 decode_vector for mul_mat_id B matrix loads (llama/23991)

This allows vec4 loads of the B elements. Also increase BK to 64 when this is
enabled. Neither of these alone is consistently faster, but together these give
a nice speedup.

In ggml-vulkan.cpp, we need to make sure the B matrix alignment and stride are
multiples of 4.

cuda: reset cuda context after reading memory size (llama/23935)

* cuda: reset device in get_memory function if no backend is active

* also count device and host buffers

* exclude hip and musa from counting and device reset

* use device mutex instead of atomic

* undo backend_free function move

ci : pin github actions to commit sha's (#3865)

coreml : fix --quantize crash for mlprogram format; fix --optimize-ane label (#3868)

commit 8b92060 switched ct.convert() to mlprogram, but did not update
the --quantize path.  quantize_weights() from
neural_network.quantization_utils only works with the legacy
neuralnetwork format.  Running with --quantize crashed with:

  Exception: MLModel of type mlProgram cannot be loaded just from the
  model spec object. It also needs the path to the weights file.

Fix: pass compute_precision=ct.precision.FLOAT16 into ct.convert() when
--quantize is set.  This matches the original intent of nbits=16 (F16
storage) without changing the quantization scheme or model accuracy.

Also fix the three boolean CLI flags (--encoder-only, --quantize,
--optimize-ane) to use a _str_to_bool helper so that both
  --flag True
and
  --flag False
parse correctly.  The type=bool form accepted "False" as True because
bool("False") == True.

Remove the "currently broken" label from --optimize-ane: the ANE path
(WhisperANE with Conv2d attention and LayerNormANE) converts and loads
correctly with both PyTorch 2.x and coremltools 9.x.

talk-llama : sync llama.cpp

sync : ggml

ggml : bump version to 0.14.0 (ggml/1533)

sync : ggml

HIP: add gfx1152 and gfx1153 to RDNA3.5 (llama/24129)

metal : fix im2col 1D case (audio models) (llama/24220)

vulkan: check coopmat2 features before reporting support (llama/24186)

opencl: improve get_rows, cpy, concat and q6_k flat gemv (llama/24160)

* opencl: allow multiple workgroups for large rows

* opencl: improve small cpy

* opencl: packed concat for small input

* opencl: tweak flat q6_K gemv, increase N_DST and remap threads

vulkan: add fwht support for Intel with shmem reduction (llama/23964)

* vulkan: add fwht support for Intel with shmem reduction

* don't use N as workgroup size

* disable subgroup shuffle on MoltenVK AMD

* disable fwht shader on Intel Windows due to driver bug

kleidiai : dynamic chunck-based scheduling for hybrid execution (llama/23819)

CUDA: enroll mul_mat_vec_q_moe into pdl (llama/24087)

* Enroll mul_mat_vec_q_moe into PDL, boosting MTP performance on BW

Data collected on a B4500:

Before
```
(llama.cpp) ➜  llama.cpp git:(master) ✗ python mtp-bench.py
  code_python        pred= 192 draft= 150 acc= 116 rate=0.773 tok/s=202.8
  code_cpp           pred= 192 draft= 147 acc= 117 rate=0.796 tok/s=212.8
  explain_concept    pred= 192 draft= 161 acc= 110 rate=0.683 tok/s=196.4
  summarize          pred= 192 draft= 138 acc= 122 rate=0.884 tok/s=226.6
  qa_factual         pred= 192 draft= 138 acc= 121 rate=0.877 tok/s=225.1
  translation        pred= 192 draft= 158 acc= 112 rate=0.709 tok/s=201.5
  creative_short     pred= 192 draft= 160 acc= 110 rate=0.688 tok/s=197.2
  stepwise_math      pred= 192 draft= 150 acc= 115 rate=0.767 tok/s=209.2
  long_code_review   pred= 192 draft= 148 acc= 116 rate=0.784 tok/s=208.9
```
After
```
(llama.cpp) ➜  llama.cpp git:(master) ✗ python mtp-bench.py
  code_python        pred= 192 draft= 150 acc= 116 rate=0.773 tok/s=211.9
  code_cpp           pred= 192 draft= 147 acc= 117 rate=0.796 tok/s=224.6
  explain_concept    pred= 192 draft= 161 acc= 110 rate=0.683 tok/s=207.8
  summarize          pred= 192 draft= 138 acc= 122 rate=0.884 tok/s=240.2
  qa_factual         pred= 192 draft= 138 acc= 121 rate=0.877 tok/s=238.5
  translation        pred= 192 draft= 158 acc= 112 rate=0.709 tok/s=213.4
  creative_short     pred= 192 draft= 160 acc= 110 rate=0.688 tok/s=208.8
  stepwise_math      pred= 192 draft= 150 acc= 115 rate=0.767 tok/s=221.7
  long_code_review   pred= 192 draft= 148 acc= 116 rate=0.784 tok/s=220.7
```

Server launched with:
```
➜  llama.cpp git:(osimons/enroll_mul_mat_vec_q_moe_into_PDL) ✗ ./build-x64-linux-gcc-reldbg/bin/llama-server \
    -m /mnt/share/gguf/unsloth/Qwen3.6-35B-A3B-MTP-GGUF/Qwen3.6-35B-A3B-UD-Q4_K_M.gguf -dio \
    --spec-type draft-mtp \
    --spec-draft-n-max 2 \
    -ngl all \
    -fa on \
    --host 0.0.0.0 \
    --port 8080 -np 1 --chat-template-kwargs "{\"preserve_thinking\": true}"
```

* LC to overlap with following kernels

sycl : port multi-column MMVQ from CUDA backend (llama/21845)

mmvq:

Port the ncols_dst optimization from ggml-cuda/mmvq.cu to SYCL.
Read weights once per dispatch instead of once per column.
Covers all standard quant types + reorder paths for Q4_0, Q8_0,
Q3_K, Q4_K, Q5_K, Q6_K. IQ types (except IQ4_XS) excluded due to
incompatible vec_dot signatures.

ggml-sycl:

The weight reorder was only bootstrapped on single-token mat-vec
(ne[1] == 1). Speculative / MTP verify issues only multi-column mat-vec,
so it never triggered the reorder and ran on the slower non-reorder
kernel. Bootstrap it on small multi-column batches (ne[1] <= 8) too.

ggml: vectorize ggml_vec_dot_q4_1_q8_1 with WASM SIMD128 (llama/22209)

* ggml: vectorize ggml_vec_dot_q4_1_q8_1 with WASM SIMD128

Optimize the inner loop of ggml_vec_dot_q4_1_q8_1_generic using
WASM SIMD128 intrinsics, gated behind #ifdef __wasm_simd128__ so
non-wasm builds are completely unaffected.

Approach:
- single wasm_v128_load covers all 32 packed 4-bit weights
- nibbles unpacked via AND/SHR into two u8x16 registers
- widened to i16 before multiply (WASM SIMD has no i8*i8 instruction)
- 4x wasm_i32x4_dot_i16x8 calls accumulate all 32 element pairs
- horizontal reduce via 4x wasm_i32x4_extract_lane

Benchmark (node v25, emcc -O3 -msimd128, 64 blocks x QK8_1=32,
200k iterations):

| impl   | ns/call | speedup |
|--------|---------|---------|
| scalar |   880.7 |   1.00x |
| simd   |   257.8 |   3.42x |

Correctness verified against scalar reference across 10 random seeds
with exact output match.

* ggml: move q4_1_q8_1 WASM SIMD implementation to wasm backend

Relocate the SIMD128 implementation of ggml_vec_dot_q4_1_q8_1 to ggml/src/ggml-cpu/arch/wasm/quants.c to follow architecture-specific layout. Restore the generic implementation in ggml/src/ggml-cpu/quants.c.
Move for loop in the else block.

* ggml: use generic q4_1_q8_1 fallback in wasm backend

metal : reduce rset heartbeat from 500ms -> 5ms (llama/24074)

ggml-webgpu: FlashAttention refactor + standardize quantization support (llama/23834)

* Start work on flash_attn refactor

* Refactor

* Split k/v quantization

* Refactor and abstract quantization logic for flash_attn and mul_mat

* Add quantization support to tile path

* formatting

* Move to functions, add a check

ggml-cpu: extend RVV quantization vec dot to higher VLENs (llama/22754)

* ggml-cpu: add rvv 512b,1024b impls for iq4_xs

* ggml-cpu: refactor; add rvv 512b, 1024b impls for q6_K, i-quants

* ggml-cpu: refactor; add 512 and 1024 implementations of tq3_s, iq3_xxs, iq2_s, iq2_xs, iq2_xxs

improve iq2_xs impl for rvv 256

Co-authored-by: Rehan Qasim <redacted>
---------

Co-authored-by: taimur-10x <redacted>
Co-authored-by: Rehan Qasim <redacted>

Avoid PDL race conditions by disabling __restrict__ when PDL is used (llama/24030)

* Removes __restrict__ from PDL kernel headers due to incompatibility with
PDL. Adds preprocessor directives based on arch in kernel body to add
__restrict__ to retain performance on older architectures.

* Simplifies new __restrict__ usage via macro

* Add hopper to PDL __restrict__ fix.

Co-authored-by: Oliver Simons <redacted>
---------

Co-authored-by: Oliver Simons <redacted>

ggml-cpu: use runtime SVE width in FWHT (llama/24059)

cuda: reserve space for quantize kv-cache at startup (llama/23907)

* cuda: reserve space for quantize kv-cache at startup

* address review comments

* remove forward decl

Co-authored-by: Johannes Gäßler <redacted>
* remove assert in ggml-cuda.cu

Co-authored-by: Johannes Gäßler <redacted>
---------

Co-authored-by: Johannes Gäßler <redacted>

opencl: use flat variants of q4_K and q6_K gemv for very large M (llama/24006)

hexagon: profiler output fix and script updates (llama/24042)

* hex-ops: fix profiler output (ie remove the redundant NONEs)

* hex-prof: update profiling script to support tot.usec column

hexagon: MUL_MAT, MUL_MAT_ID, FLASH_ATTN and GDN cleanup and optimizations for latest models (llama/23989)

* hex-mm: initial support for F32 * F32 -> F32 matmuls

* hex-rms-norm: fix src1 stride use in fused rms_norm_mul

* hex-ops: clear spad pointers in the ops that clober it

This fixes an odd case where fused rms-norm-mul was failing but only in qwen3.5-2B and only at searth op-bath sizes.

* hmx-mm: add support for F32 * F32 -> F32 matmul_2d on HMX

Decided to use Q4_0 * F32 -> F32 matmul for this.
Q4_0 gets dequantized and tiled into F16, and here we quantize and tile F32 into F16.
Super simple and pretty efficient.

* hmx-mm: route f16 2D matmuls through the same kernel used for all other types

* hmx-mm: re-introduce pipelined vs non-pipelined mode that we used to have but is much more generic way

This update futher improves matmul performance and at the same time removes most of the redudant logic
we had in different paths.

* hmx-fa: slighlty improved pipeline simimar to matmul updates

* hmx-mm: initial version of MAT_MUL_ID support for HMX

* hmx-mm: fixed mxfp4 handling for MUL_MAT_ID

* hex-gdn: optimize GATED_DELTA_NET

DMA prefetch/double-buff, vectorize everything with HVX, in other words -- the usual :)

* hmx-mm: missed one more case where we can use fastmod

* hexagon: update DCVS settings for a slight perf bump

* hmx-fa: use fastdiv in hmx-flash-attn

* hmx-fa: precompute slope values to avoid disrupting the inner loop

* hvx-utils/fa: new HVX helpers for powf and logf and using those to speed up FA alibi

* hex-ops: fixed a bug in fusion logic that was messing up the order of the src tensors when some srcs are empty

* hex-fa: correctly fallback to HVX if we have sinks or the dims are not quite right

hexagon: add gelu_quick (llama/24007)

clean up unused variables warnings (llama/23975)

opencl: fix compiler warnings for non-adreno path (llama/23922)

* opencl: fix compiler warnings for non-adreno path

* opencl: fix const cast warning

revert to using global_invocation_id for cpy shader (llama/23955)

opencl: add basic support for q5_0 and q5_1 (llama/23548)

* opencl: add general q5_0 support

* opencl: add general q5_1 support

* opencl: support non-uniform workgrp size

---------

Co-authored-by: Li He <redacted>

metal: template GLU kernels to support f16/f32 (llama/23882)

Drops the hardcoded f32 GLU kernels in favor of a single template. We now load/store in the native tensor type (half or float) to save memory bandwidth, but keep the actual ALU compute in float to avoid exploding math in geglu/swiglu. Also opened up the dispatch gate to allow f16 inputs.

vulkan: don't hold the device mutex while compiling pipelines (llama/23641)

* vulkan: don't hold the device mutex while compiling pipelines

We need to hold a lock while we traverse all pipelines and lazily initialize
them, but we don't need to hold it while the pipeline is being compiled. And
it doesn't need to be the same lock as the device mutex. We call load_shaders
each time a pipeline is needed, so we only need to compile that one pipeline
(and, for example, don't want to end up compiling a pipeline that another
thread should be compiling).

* remove 'needed'

vulkan: reduce host memory lock contention (llama/23376)

* vulkan: reduces lock contention

* replace unique_lock with lock_guard

TP: quantized KV cache support (llama/23792)

* TP: quantized KV cache support

* fix partial view

* remove overly strict assert

vulkan: Block-load Q3_K/Q6_K block data and subtract on 32b ints (llama/23056)

Q2_K/Q3_K/Q6_K do much better when using MMVQ on Intel BMG even
though they're only 2-byte aligned, and Q3_K still wins on
NVIDIA as well.

mesa isn't all that great at coalescing back-to-back loads from
alternating arrays, so we force it instead. Further, we can do
subtraction directly on a full int32_t rather than an i8vec4
with bit twiddling because the high bit is always free to start.

On Intel BMG on mesa, the switch to MMVQ provides an immediate
~57% perf increase in tg128 for unsloth/Qwen3.5-9B-GGUF:Q3_K and
~78% perf increase in tg128 for unsloth/Qwen3.5-9B-GGUF:Q6_K.

The futher switch to block loads leads to a ~24% perf increase in
tg128 for unsloth/Qwen3.5-9B-GGUF:Q3_K and a ~48% perf increase in
tg128 for unsloth/Qwen3.5-9B-GGUF:Q6_K.

Finally, Xe2 wins on MMVQ even for small k, so we take the NVIDIA
override for K quants on Xe2 as well.

vulkan: Removed unused functions (llama/23175)

Support Q4_1, Q5_0, Q5_1 in Flash-attention (llama/23812)

* support Q4_1, Q5_0, Q5_1

* update ut case

Add more types in GET_ROWS OP (llama/23710)

* add to support Q1_0, NVFP4, IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ1_S, IQ1_M, IQ3_S, IQ4_NL, IQ4_XS, I32, MXFP4, Q2_K, Q3_K, Q5_K, and Q6_K in GET_ROWS OP

* correct the link

sycl : Optimize Q3_K mul_mat by reorder (llama/23725)

opencl: support bf16 by converting to f16 (llama/23839)

metal : restore im2col implementation for large kernels (llama/23901)

ggml : add some lsx support (llama/23798)

* loongarch : optimize LSX fp16 load/store with native intrinsics

Use __lsx_vfcvtl_s_h and __lsx_vfcvt_h_s instead of scalar loops in
__lsx_f16x4_load and __lsx_f16x4_store.

* loongarch : add LSX implementation for q8_0 dot product

* loongarch : add LSX implementation for q6_K dot product

* loongarch : add LSX implementation for iq4_xs dot product

* Improve reduce ops when sun int16 pairs to int32

vulkan: add Flash Attention support for BFloat16 KV cache (llama/23420)

* vulkan: add flash attention bf16 kv support

* vulkan: bf16 FA coopmat1 support

* vulkan: bf16 FA coopmat2 support

* fix FA bf16 f32 fallback

* fix FA bf16 coopmat1 shader

* fix FA bf16 coopmat2 shader

* code cleanup

* cleanup comment change

* address feedback

* add O_TYPE for cm2 FA

* use O_TYPE for gqaStore function

* reduce BFLOAT16 ifdefs

ggml-webgpu: Check earlier for WebGPU required features (llama/23879)

ggml-webgpu: add q4_0/q8_0 SET_ROWS (llama/23760)

* Add q8_0 and q4_0 set_rows

* Add fast(er) quantization set_rows path

* formatting/naming

* a little more naming

* Remove unused constant

* Don't override other override

* Avoid bitcast

* Narrow relaxation

CUDA: Check PTX version on host side to guard PDL dispatch (llama/23530)

* CUDA: Check PTX version on host side to guard PDL dispatch

Checking on `__CUDA_ARCH_LIST__` alone is insufficient for JIT, as this
variable doesn't differentiate between compiling for say sm_90, sm_90a
or sm_90f (so forward-jittable PTX vs. arch/family-specific PTX).

Thus, one can have a bug when compiling with
`DCMAKE_CUDA_ARCHITECTURES="89;90a"`, where current code would wrongly
dispatch to PDL on sm_90/sm_120 in forward-JIT mode.

This PR fixes this issue by checking `cudaFuncAttributes::ptxVersion` of
the incoming kernel at runtime. A check on ptxVersion alone is
sufficient, as device-codes will always be >= ptxVersion (and any
violation of this would be a severe bug in CUDA/nvcc), see:
 https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/#gpu-code-code-code

* Implement MurmurHash3 mixer for better hash distribution

Magic constants were taken from boost:
https://github.com/boostorg/container_hash/blob/2698b43803c012601e6bb1a6116e83767b97986c/include/boost/container_hash/detail/hash_mix.hpp#L19-L65

* Update ggml/src/ggml-cuda/common.cuh

Co-authored-by: Johannes Gäßler <redacted>
* Address review comments, make seed non-zero

* Apply code-formatting

* Replace std::size_t -> size_t for consistency

---------

Co-authored-by: Johannes Gäßler <redacted>

model : support for DeepseekV32ForCausalLM with generic DeepSeek Sparse Attention (DSA) implementation (llama/23346)

* llama : support DeepSeek V3.2 model family (with DSA lightning indexer)

* convert : handle DeepseekV32ForCausalLM architecture

* ggml : support for f16 GGML_OP_FILL

* memory : separate hparams argument in llama_kv_cache constructor

* memory : add llama_kv_cache_dsa memory (KV cache + lightning indexer cache)

* llama : support for LLM_ARCH_DEEPSEEK32

* model : llama_model_deepseek32 implementation

* model : merge two scale operations into one in DSA lightning indexer implementation

* chore : remove unused code

* model : support NVFP4 in DeepSeek V3.2

Co-authored-by: Sigbjørn Skjæret <redacted>
* memory : refactoring TODO

Co-authored-by: ggerganov <redacted>
---------

Co-authored-by: Stanisław Szymczyk <redacted>
Co-authored-by: Sigbjørn Skjæret <redacted>
Co-authored-by: ggerganov <redacted>

ci : add ccache to build-sycl [no ci] (#3859)

ci : add HF_TOKEN to docker.yml workflow [no ci] (#3861)

This commit adds the HF_TOKEN secret to the docker workflows to avoid HF
rate limiting which currently sometimes causes the jobs to fail.

Refs: https://github.com/ggml-org/whisper.cpp/actions/runs/27053852601/job/79854251771

ci : add ccache to quantize, vad, and wasm jobs (#3860)

* ci : add ccache to build-quantize

* ci : add ccache to build-vad

* ci : add ccache to build-wasm [no ci]

ci:  build-windows action slimming (#3858)

* ci : remove base-devel and git from msys2 job

This commit removes the above packages as they might not be required and
could help reduce the github cache size.

* ci : try reducing the installs to only the compilers

ci : use emscripten-core and pin version (#3857)

This commit updates the setup emscripten sdk jobs to use emscripten-core
instead of mymindstorm and also pins the commit sha for the version
instead of using a version tag.

ci : pin github actions to commit SHAs (#3856)

This commit pins github actions used to the same commi SHAs that
llama.cpp uses.

ci : use ccache instead of sccache for windows-cublas [no ci] (#3855)

This commit updates the Install cache step to use ggml-org/ccache-action
and switched to use ccache instead of sccache.

The motivation for switching to ccache is that this is what llama.cpp
does and also there is an issue with later version of sscache:
```console

    sccache C:\PROGRA~1\NVIDIA~1\CUDA\v\bin\nvcc.exe -forward-unknown-to-host-compiler -DGGML_BACKEND_BUILD -DGGML_BACKEND_SHARED -DGGML_CUDA_PEER_MAX_BATCH_SIZE=128 -DGGML_SCHED_MAX_COPIES=4 -DGGML_SHARED -D_CRT_SECURE_NO_WARNINGS -D_XOPEN_SOURCE=600 -Dggml_cuda_EXPORTS -DCMAKE_INTDIR=\"Release\" -ID:\a\whisper.cpp\whisper.cpp\ggml\src\ggml-cuda\.. -ID:\a\whisper.cpp\whisper.cpp\ggml\src\..\include -isystem "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v\include" -Xcompiler="-MD -O2 -Ob2" -DNDEBUG -std=c++17 -arch=native -use_fast_math -extended-lambda -Xcompiler /Zc:preprocessor -MD -MT ggml\src\ggml-cuda\CMakeFiles\ggml-cuda.dir\Release\allreduce.cu.obj -MF ggml\src\ggml-cuda\CMakeFiles\ggml-cuda.dir\Release\allreduce.cu.obj.d -x cu -c D:\a\whisper.cpp\whisper.cpp\ggml\src\ggml-cuda\allreduce.cu -o ggml\src\ggml-cuda\CMakeFiles\ggml-cuda.dir\Release\allreduce.cu.obj -Xcompiler=-Fdggml\src\ggml-cuda\CMakeFiles\ggml-cuda.dir\Release\,-FS

    sccache: encountered fatal error

    sccache: error: Could not parse shell line

    sccache: caused by: Could not parse shell line

    ```
```

ci : only publish/push docker images daily (#3854)

This commit updates the docker workflow to be triggered on a schedule or
manually.

ci : refactor + optimize (#3847)

* ci : add ccache clear action

* ci : split self-hosted GPU jobs into build-self-hosted.yml

Extract self-hosted runner jobs from build.yml into a dedicated
build-self-hosted.yml following the llama.cpp pattern:
  - gpu-cuda (NVIDIA Linux)
  - gpu-vulkan-nvidia-cm (NVIDIA Linux)
  - gpu-vulkan-nvidia-cm2 (NVIDIA Linux + COOPMAT2)
  - gpu-metal (macOS ARM64)
  - gpu-vulkan (macOS ARM64)

GitHub-hosted CPU jobs remain in build.yml.

Assisted-by: llama.cpp:local pi
* ci : split release jobs into release.yml

Extract release-related jobs from build.yml into a dedicated
release.yml following the llama.cpp pattern:
  - determine-tag
  - windows (Win32/x64, SDL2)
  - windows-blas (Win32/x64, OpenBLAS)
  - windows-cublas (x64, CUDA 11.8/12.4)
  - ios-xcode-build
  - bindings-java (depends on windows)
  - release (artifact aggregation + GitHub release)

CoreML job stays in build.yml with its own local tag calculation.

Assisted-by: llama.cpp:local pi
* ci : remove bindings-java job from release.yml

Assisted-by: llama.cpp:local pi
* cont : add manual trigger for build.yml

* cont : remove obsolete ifs

* ci : extract sanitizer job to bild-sanitize.yml

* ci : extract linux jobs into build-linux.yml

* ci : extract macos jobs to build-macos.yml

* ci : extract gcc jobs to build-gcc.yml

* ci : extract clang jobs to build-clang.yml

* ci : extract sycl jobs to build-sycl.yml

* ci : extract windows jobs to build-windows.yml

* ci : extract emscripten job to build-wasm.yml

* ci : extract android jobs into build-android.yml

* ci : extract quantize job to quantize.yml

* ci : extract coreml job into coreml.yml

* ci : extract vad job to vad.yml

* ci : extract cpu jobs to build-cpu.yml

* ci : make naming of yml files consistent

* ci : add --fail to curl download and propagate

This commit adds the --fail option to the model download scripts so that
if the model download returns a server error this is picked up. This is
then detected in run.sh and a error message is displayed and the script
stops and returns an error.

The motivation for this is that currently it is possible for the model
download to fail but this script proceeds and instead of a model file
the contents will be an html page probably with the error. This will
then cause the model to not be able to load due to a missing magic
number. I'm not sure we can do much about the downloading failing,
perhaps a retry but at least this will give a clearer error message.

Refs: https://github.com/danbev/whisper.cpp/actions/runs/26866349389/job/79230794512

* ci : enable command traces to see download command in use

* ci : add retry functionality to download model script

This commit adds curl retry options to the model download script.

The motivation is that currently when CI jobs run huggingface rate limit
the requests and return:
```console
curl: (22) The requested URL returned error: 429
```
This is an attempt to work around this and if it does not work then we
can an authorization token.

* ci : extract freebsd job to build-freebsd.yml

This job has been commented out as it has been flaky in the past. I'll
monitor this and if it continues to be unreliable we can disable it in
the github actions GUI instead of commenting it out like we did before.

* ci : add ccache to jobs (non-docker builds)

The ccache will only be saved on pushed to master.

* ci : bump ccache-action version to v1.2.21

The motivation for this is that the save parameter does not seem to work
with the current version.

* ci : add ccache to docker jobs in build-linux.yml

* ci : add debug statements to linux docker build

* ci : set CCACHE_DIR for build-linux.yml

* ci : add ccache to the remaining docker jobs

* ci : remove build-linux.yml

This commit remove build-linux.yml as the same jobs are also run by
build-gcc.yml, with the exception that build-gcc.yml also run ctest).
So keeping build-gcc.yml and removing the redundant build-linux.yml.

* ci : add linux build artifacts to release

* ci : revert to hendrikmuhs/ccache-action for win job

This is currently causing the following failure:
```console
sccache C:\PROGRA~1\NVIDIA~1\CUDA\v\bin\nvcc.exe -forward-unknown-to-host-compiler -DGGML_BACKEND_BUILD -DGGML_BACKEND_SHARED -DGGML_CUDA_PEER_MAX_BATCH_SIZE=128 -DGGML_SCHED_MAX_COPIES=4 -DGGML_SHARED -D_CRT_SECURE_NO_WARNINGS -D_XOPEN_SOURCE=600 -Dggml_cuda_EXPORTS -DCMAKE_INTDIR=\"Release\" -ID:\a\whisper.cpp\whisper.cpp\ggml\src\ggml-cuda\.. -ID:\a\whisper.cpp\whisper.cpp\ggml\src\..\include -isystem "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v\include" -Xcompiler="-MD -O2 -Ob2" -DNDEBUG -std=c++17 -arch=native -use_fast_math -extended-lambda -Xcompiler /Zc:preprocessor -MD -MT ggml\src\ggml-cuda\CMakeFiles\ggml-cuda.dir\Release\allreduce.cu.obj -MF ggml\src\ggml-cuda\CMakeFiles\ggml-cuda.dir\Release\allreduce.cu.obj.d -x cu -c D:\a\whisper.cpp\whisper.cpp\ggml\src\ggml-cuda\allreduce.cu -o ggml\src\ggml-cuda\CMakeFiles\ggml-cuda.dir\Release\allreduce.cu.obj -Xcompiler=-Fdggml\src\ggml-cuda\CMakeFiles\ggml-cuda.dir\Release\,-FS
sccache: encountered fatal error
sccache: error: Could not parse shell line
sccache: caused by: Could not parse shell line
```

Refs: https://github.com/danbev/whisper.cpp/actions/runs/26883673904/job/79290017353

* ci : make static linux artifacts

* ci : make linux release artifact names consistent

This commit removes the tag form the linux release artifacts to be
consistent with the existing artifacts.

If we want to include the tag then we can do that in a follow-up PR.

* ci : fix linux zip files to have a directory

* ci : add HF_TOKEN secret for HF download authorization

This is to avoid the HR rate limiting when downloading model.

---------

Co-authored-by: Daniel Bevenius <redacted>

whisper : catch C++ exceptions in whisper_init_with_params_no_state (#3831)

whisper_model_load() can throw instead of returning false: std::runtime_error
from this file (failed ggml context / no compatible buffer type), or
vk::SystemError / vk::OutOfDeviceMemoryError from the ggml-vulkan backend during
device/buffer allocation.

whisper_init_* are extern "C", so a C++ exception unwinding across that boundary
aborts non-C++ callers (Rust via whisper-rs, Go via cgo) -- on Windows
STATUS_STACK_BUFFER_OVERRUN (0xC0000409) -- even though the function already
returns NULL on failure. Wrap whisper_model_load() in try/catch and route any
throw into the existing NULL-return path.

Co-authored-by: Claude Opus 4.7 <redacted>

server : merge split utf-8 token text in verbose json (#3850)

cmake : do not assume /usr/lib library installation. (#3693)

Current `pkgconfig` configuration file installation path and its
contents assume libraries are installed under `/usr/lib` and this is not
always the case, for instance `/usr/lib64` is quite possible under
Gentoo Linux.

Thus use the `CMAKE_INSTALL_LIBDIR` variable instead of a hardcoded
`lib`.

release : v1.8.6

ci : fix path to whisper.h in examples.yml [no ci] (#3842)

This commit updates the include path to whisper.h and also ensures that
this is only built on pushes to master.

ci : fix self-hosted paths to mnt

pi : add config

[no ci]

ci : remove obsolete self-hosted label

common : pass sample rate to `ffmpeg_decode_audio()`

common : re-implement `ffmpeg-transcode.cpp` + clarify ffmpeg usage (#3846)

* examples : remove ffmpeg-transcode.cpp

* examples : implement ffmpeg-transcode.cpp

Assisted-by: llama.cpp:local pi
* common : switch from WHISPER_FFMPEG -> WHISPER_COMMON_FFMPEG

sync : ggml

ggml : bump version to 0.13.1 (ggml/1523)

talk-llama : sync llama.cpp

sync : ggml

cuda : disables launch_fattn PDL enrollment due to compiler bug (llama/23825)

meta : Add missing `buffer` set in allreduce fallback !COMPUTE clear (llama/23480)

Without this at least the vulkan backend will skip the `* 0` for
!COMPUTE tensors, causing corrupt output.

hexagon: basic/generic op fusion support and RMS_NORM+MUL fusion (llama/23835)

Updating infra to enable op fusion and using RMS_NORM+MUL as the use-case.

opencl: move backend info printing into its own function (llama/23702)

* opencl: move backend info print into its own function

* opencl: move new log line

* opencl: fix for non adreno path

ggml: auto apply iGPU flag CUDA/HIP if integrated device (llama/23007)

mmvq Optim: add MMVQ_PARAMETERS_TURING(mmvq_parameter_table_id) for … (#23729)

* mmvq Optim:  add MMVQ_PARAMETERS_TURING(mmvq_parameter_table_id) for SM75 TURING

* avoid a mismatch for JIT compilation of Turing device code for Ampere or newer

Co-authored-by: Johannes Gäßler <redacted>
---------

Co-authored-by: Copilot <redacted>
Co-authored-by: Johannes Gäßler <redacted>

CUDA: route batch>=4 quantized matmul to MMQ on AMD MFMA hardware (llama/23227)

* CUDA: per-quant MMVQ/MMQ batch threshold on AMD MFMA hardware

The dispatcher uses a single global threshold (MMVQ_MAX_BATCH_SIZE = 8)
to choose between mul_mat_vec_q (per-row GEMV) and mul_mat_q (MFMA-tiled
GEMM) for quantized matmul. On AMD CDNA, the optimal crossover differs
substantially by quant family because the per-row GEMV cost is dominated
by dequantisation, not the dot-product itself: K-quants pay a heavier
super-block decode and so MMQ wins sooner; legacy and IQ quants have
lean decode and stay ahead until the batch fully populates an MFMA tile.

This patch introduces ggml_cuda_should_use_mmvq(type, cc, ne11) -> bool,
mirroring the existing ggml_cuda_should_use_mmq, and gates per-quant
thresholds on amd_mfma_available(cc):

  Q3_K, Q4_K, Q5_K  : MMVQ <= 3   (MMQ wins from batch=4: +5% .. +76%)
  Q2_K, Q6_K        : MMVQ <= 5   (MMQ wins from batch=6: +8% .. +35%)
  others            : MMVQ <= 8   (legacy & IQ regress under MMQ; unchanged)

Non-AMD-MFMA paths (NVIDIA, RDNA, CDNA1 without MFMA) are byte-identical
to master. GGML_CUDA_FORCE_MMVQ=1 restores the original global threshold
for A/B testing.

Measured on MI250X (gfx90a, ROCm 7.2.1) with Llama-3.2-3B-Instruct,
llama-bench pp512 across all 20 supported quants, ubatch 1..8, 10 reps.
Full table in PR description.

  Selected pp512 throughput (tok/s, ub=8):
    Q4_K_S:  559 -> 940  (+68%)
    Q5_K_S:  503 -> 884  (+76%)
    Q3_K_S:  629 -> 879  (+40%)
    Q2_K  :  615 -> 809  (+32%)
    Q6_K  :  582 -> 776  (+33%)

  Selected pp512 throughput (tok/s, ub=4):
    Q4_K_S:  444 -> 480  (+ 8%)
    Q4_0  :  682 -> 685  (+ 0%)   (no regression - retains MMVQ)
    IQ4_XS:  706 -> 698  (- 1%)   (no regression - retains MMVQ)

* CUDA: address review — inline MMVQ batch table, drop env hatch & doc block

* tune kernel selection logic for CDNA1

---------

Co-authored-by: Johannes Gäßler <redacted>

hexagon: minor refresh for HMX FA and MM (llama/23796)

* hex-fa: clean up qf32/fp32 handling and stride handling

* hex-fa: fix corner case fp NAN issues that were cause bad output from gemma4 on v79

* hex-fa: vectorize leftover handling

* hex-fa: avoid HVX fallback during token gen HMX has more FP16 compute capacity

* hmx-mm: remove dead code

* hmx-mm: use fastdiv in x4x2 dequant

* hmx-mm: sandwich dequant and scatter to improve perf

* hmx-mm: fixed rebase conflicts

* hmx-mm: further improve weight dequant by doing early type dispatch and precomputing fastdiv

* hmx-mm: an even earlier dispatch for per-type dequant

* hmx-mm: dequant linear types like q4_0 and q4_1 without the LUTs

This is a bit faster than LUT.

* hex-cmake: one more tweak for lto

---------

Co-authored-by: Trivikram Reddy <redacted>

vulkan: fast path for walsh-hadamard transform (llama/23687)

* vulkan: fast path for walsh-hadamard transform

* disable for intel due to segfault

vulkan: fix wrong index variable in inner loop (llama/23665)

vulkan: Fix memory logger unsafe iterator access (llama/23667)

cuda : fix KQ mask offset integer overflow in fattn MMA kernel (llama/23610)

Co-authored-by: Stanisław Szymczyk <redacted>

ggml: fixed Arm SVE usage bug in vec.h, vec.cpp (llama/22841)

* Updated vec.h/vec.cpp code to accumulate to F32 rather than F16

Change-Id: I0cb789347f2bf60ffaf9047319f727e788c825f8

Signed-off-by: Martin Klacer <redacted>
Co-authored-by: Milos Puzovic <redacted>