opencl: flush profiling batch at shutdown for incomplete batches (llama/25016)

CUDA: Various fixes to `cpy.cu` (llama/25000)

* Add failing test-case to test-backend-ops

Extracted from https://github.com/ggml-org/llama.cpp/issues/24072

* Minimize repro with help of AI

N = 8 * (65535 - 1) + 1 = 524273

* Port and adjust workaround from https://github.com/LostRuins/koboldcpp/commit/0ba798341e0c70517cb226cb63c966b086a3b5b3

Fall-back should share code, also relax y-z constraint to be inclusive

* Add test-case + fallback also for y dim

* Fix x-guards which is 2^{31}-1, so inlusive of INT_MAX

* Fix overflow problems for transposed copy kernel

ggml : address integer overflows in binary ops CUDA implementation (llama/24706)

* ggml : address integer overflows in binary ops CUDA implementation

* ggml : add size_t casts to avoid integer overflows

* ggml : add more asserts checking integer overflows in binary ops CUDA implementation

---------

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

sycl : support --split-mode tensor (llama/24152)

* Sycl tp stage1 (llama/1)

* SYCL: tensor parallelism (--split-mode tensor) for dual-GPU

Adds the comm_init/comm_free/comm_allreduce_tensor trio that the
meta-backend queries via get_proc_address to enable backend-specific
all-reduce, mirroring the pattern used by ggml-cuda.cu.

For N=2 (the common dual-GPU case) implements a degenerate ring
all-reduce with two size-branched paths:

  * Small (nelem < 32768): FP32 direct memcpy + per-device ADD kernel
    chained via depends_on(memcpy_event). 4 SYCL submissions/call.

  * Large (nelem >= 32768): BF16-compressed. Each device compresses
    FP32 -> BF16 in a local outbox, cross-device memcpys to the peer's
    inbox (HALF the PCIe bytes), then decompresses + adds into the
    local FP32 partial. 6 SYCL submissions/call but PCIe bytes halved
    -- wins for any tensor where PCIe dominates kernel time.

Threshold and BF16 path pattern mirror the CUDA NCCL allreduce.

Storage: ONE persistent uint8_t buffer per device, 4 * nelem bytes
(matches both path layouts: FP32 nelem floats; BF16 outbox+inbox =
2 * nelem uint16_t each). Single alloc+free per device keeps the
SYCL pool's strict-LIFO invariant trivial.

Initial impl handles N=2 FP32 contiguous tensors. Other cases return
false, causing the meta-backend to use its generic butterfly fallback.

Per-call sync is intentionally omitted. SYCL in-order queue semantics
ensure that the meta-backend's next compute on the same per-device
queue waits for our final ADD, and the next allreduce's first op on
the same persistent buffer waits via the same queue. Only comm_free
does an explicit final wait.

OneCCL is NOT used: OneCCL 2021.17 hardcodes single-device-per-process
in communicator_impl.hpp:47 (condition devices.size() == 1), which is
incompatible with llama.cpp's single-process multi-GPU model.

Measured on dual Intel Arc Pro B70 (NEO 26.05.x, oneAPI 2025.3 +
DPC++ nightly):

  Llama-3.3-70B Q4_K_M, -sm tensor -fa 1 -ctk f16 -ctv f16:
    pp512 = 377.08 t/s  (vs 313.65 layer mode = +20.2%)
    tg128 = 17.40 t/s   (vs   9.74 layer mode = +78.6%)

  Qwen3-Coder-Next-80B-A3B Q3_K_M (MoE):
    pp512 = 216.56 t/s  (vs 156.58 meta-backend butterfly = +38.3%)
    tg128 = 17.60 t/s   (vs  14.31 meta-backend butterfly = +23.0%)

  Qwen3-4B Q4_K_M:
    pp64  = 984.51 t/s, tg16 = 49.29 t/s

Llama-3.3-70B in SYCL TP now comfortably beats production layer mode
on both prefill and decode. Coder-Next-80B-A3B (MoE) also wins on
both — the BF16 path is what unlocks the many-medium-allreduces
prefill pattern.

Build/CMake: no changes. No new dependencies. ~210 lines added across
ggml-sycl.h and ggml-sycl.cpp.

* Fix comments

* documentation update to address PR feedback

* Bring over my device-to-device memcpy chagnes

* move the dev2dev_memcpy calls to the upstream 7-parameter variety

* Fix a typo and remove a trailing whitespace

sycl : fix the failed UT cases of conv_3d (llama/24900)

opencl: support non-contig rows in norm (llama/24965)

hexagon: MUL_MAT and MUL_MAT_ID rework : 32x32 tiled weight repack, kernel-params, cached graphs (llama/24954)

* hex-mm: new weight layout and fusion updates

* hvx-mm: unroll the new tiled vec_dots to optimize hvx register util

* hex-mm: optimize dyn.quant format for q8_0 and q8_1 to reduce overhead in vec_dots.

* hvx-mm: parallel quantizer per block for large rows

* hvx-mm: simplify and futher optimize dyn.quant and vec_dots

* hvx-mm: keep intermediate per tile accumulators in fp16

* hmx-mm: optimize weight dequant by aligning the repacked tiles with the DMA

* hmx-mm: remove qweight scratch and just use vtcm_weight

* hmx-mm: remove all unused and obsolete code

* hmx-mm: the new tiled repack format is here to stay -- rename all x4x2 to _tiled

* hmx-mm: improve activation processing with dma prefetch

* hex-mm: fix hmx/hvx fallback logic and MUL_MAT_ID allocation (unbreaks OLMoE)

* hex-mm: align the weight tiles with dma just like we did in hmx-mm

* hex-mm: factor out common mm bits into htp/matmul-ops.h

* hex-mm: start moving mm kernel selection to the host

* hex-mm: move all of the matmul param compute into the host

* hmx-mm: restore pipelined mode

* hmx-mm: unroll the dequant functions to optimize register usage

* hmx-mm: further improve activation process

* hex-mm: use vtcm_seq_alloc for all vtcm allocations and define more common functions

* hex-mm: improve mm optimizer to acount for number of activation threads

* hex-mm: fix matmul-id kernel params selection (unbreaks OLMoE and LFM)

* hexagon: remove support for arch < v73 since HMX is now required for most use-cases

* hex-mm: cleanup naming for consistency

* hex-mm: make sure matmul fusion accounts for vtcm allocation

* hex-mm: minor cleanup for kernel_params definition

* hex-mm: replace hardcoded limits with proper checks for vtcm requirements

* hex-mm: add support for non-tiled mm as a fallback option and factor out hvx kernels into separate header

* hex-mm: remove unused functions

* hex-mm: add shorthand for MM_SELECT in run-tool script

* hvx-mm: factor out hvx/hmx microkernels and unify matmul entry and dispatch

* hex-mm: further cleanup matmul fallback path

* hex-mm: refactor matmul entry point and dispatch a bit further

* hexagon: update cmake build to enable hmx for everything

* hex-ops: optimize kernel_param updates and include summary in the logs

* hex-mm: add support for GGML_HEXAGON_MM_SELECT

* hex-mm: add hex-common header

* hex-mm: pass correct number of tasks to workpool

* hex-mm: add proper checks for no-work in dyn.quant tasks

* hex-mm: convert all quantizers into a macro

* hex-mm: fix hvx-flat fallback to pass all MUL_MAT tests

* hex-mm: vectorize q8_1 quantizer

* hex-mm: improve fused ffn mm stride handling

* hex-mm: consistent use of n_threads and pipeline in kernel_params

* hexagon: minor formatting

* hex-mm: update MUL_MAT_ID kernel_param handling to make sure host/npu are in sync

* hvx-mm: go back to accumulating in fp32 in tiled hvx kernels, more accurate and same perf

* hvx-mm: unroll the loops and remove masking that is not needed for tiled accums

* hmx-mm: optimize activation processing (slit loops, some unrolling, etc)

* hmx-mm: minor optimization for output processing

* hex-mm: consistent use of uint32_t and size_t in mm kernels

* hex-mm: remove legacy restrictions for rows to be multiple of 256

* hexagon: replace sprintf with snprintf

* hex-mm: relax hardcoded nrows checks and rely on VTCM size requirements

* hexagon: minor alignment fix

* hexagon: fix trailing spaces

* hex-mm: relax padding from 256 to 128 (leftovers)

* hex-mm: remove redundant checks for weight align to 128

we always use 2D dma for the weights and align them properly

* hmx-mm: MUL_MAT_ID better work distribution between hvx threads and hmx tracing

* hex-mm: specialize per-token mmid activation handling

* hex-profile: update python scripts to handle kernel-params section in the logging output

* hex-mm: move n_prefetch (aka dma_depth) into kernel params and remove unused fields

* hex-trace: use easier to parse format, simply and fix post-proc scripts

* hmx-mm: relax 32 row limit for output processing which helps utilization

* hmx-mm: use start-chunk idx for tracing info

* hmx-mm: parameterize activation dma pipeline

* hexagon: add support for simple graph caching to avoid recomputing kernel-params

* hex-mm: remove left-over repack functions

* hex-mm: tighten n_prefetch asserts

* hex-mm: remove duplicate round/align_up helper

* hexagon: cleanup common header used in host/npu

* hexagon: update early wakeup threshold

* hmx-mm: define cost constants and update solver to assume that repacked ne[1] is padded to 32

* hmx-mm: make precompute_matmul a bit more readable (split into smaller functions, etc)

* hex-mm: remove n_threads constraint

* hex-mm: minor formatting updates

* hex-mm: remove obsolete profiling logs

* hex-mm: restore hardcode gate to refuse lm-head to avoid repacking that tensor

vulkan: allow reducing the graph submission batches to avoid timeouts (llama/24872)

vulkan: fail the build when a shader fails to compile (llama/24450)

* vulkan-shaders-gen: fail the build when a shader fails to compile

vulkan-shaders-gen did not detect shader-compile subprocess failures, so a
broken libggml-vulkan could be produced while the build reported success and
the breakage only surfaced at run time. execute_command() discarded the child
exit code (POSIX waitpid passed nullptr for status; the Windows branch never
called GetExitCodeProcess) and string_to_spv decided success only from whether
stderr was empty, so a non-zero exit with empty stderr, or a subprocess that
failed to launch, was treated as success.

Return the child exit code from execute_command() (WEXITSTATUS on POSIX,
GetExitCodeProcess on Windows), treat a non-zero exit or non-empty stderr or a
launch exception as a failure, and record it in an atomic flag. main() checks
the flag after process_shaders() and returns EXIT_FAILURE before writing the
output files, so the build stops instead of emitting a broken backend.

Fixes #24393

Signed-off-by: liminfei-amd <redacted>
* vulkan-shaders-gen: simplify compile_failed access and drop unreachable return

Address review feedback on #24450:
- Access the std::atomic<bool> compile_failed directly (= / implicit bool)
  instead of .store()/.load(); the flag stays atomic because the worker
  threads in process_shaders() set it concurrently.
- Remove the unreachable trailing return -1 in execute_command(): on POSIX the
  child _exit()s after execvp and the parent returns (fork()<0 throws); on
  Windows the block returns the exit code.

Signed-off-by: liminfei-amd <redacted>
---------

Signed-off-by: liminfei-amd <redacted>

vulkan: Apply bias before softmax in FA, to avoid overflow (llama/24909)

vulkan: support all backend tests for SQR/SQRT/SIN/COS/CLAMP/LEAKY_RELU/NORM (llama/24582)

* vulkan: make SQR/SQRT/SIN/COS/CLAMP/LEAKY_RELU use unary.comp

* vulkan: make NORM support noncontig

* add noncontiguous row test cases for norm/l2_norm, handle this in the CPU backend and l2_norm.comp

* fix supports_op for cuda and webgpu

vulkan: Support GET_ROWS_BACK (llama/24883)

vulkan: support CONV_3D (llama/24612)

* vulkan: support CONV_3D

This is a pretty direct port of conv2d_mm.comp to CONV_3D, done by codex
and cleaned up by me.

* disable slower perf tests

vulkan: make mul_mm ALIGNED a spec constant (llama/24689)

This trims down some of the shader variant explosion and reduces binary size.

vulkan: link ggml-cpu when GGML_VULKAN_CHECK_RESULTS / RUN_TESTS are enabled (llama/24444)

The result-checking and test debug paths in ggml-vulkan.cpp call ggml_graph_compute_with_ctx() to compute a CPU reference graph, but that symbol is defined in ggml-cpu, which ggml-vulkan does not link. Enabling -DGGML_VULKAN_CHECK_RESULTS=ON (or -DGGML_VULKAN_RUN_TESTS=ON) therefore fails to link with an unresolved external (e.g. LNK2019 on MSVC, undefined reference on GCC/Clang). This regressed after ggml-cpu was split into its own library. Link ggml-cpu under those two options so the debug builds link again.

Signed-off-by: Wyatt Caldwell <redacted>

ggml-webgpu: improve MTP inference by using mat-vec path for small batches (llama/24811)

* ggml-webgpu: improve small batches decoding

* Add barrier to the NUM_COLS loop in mul-mat-vec

opencl: q8_0 gemv precision improvement (llama/24923)

support bf16 on bin_bcast OP and unary OPs (llama/24838)

* support bf16 on bin_bcast OP and unary OPs

* support the older Intel compiler than 2026.0

fix(hexagon): use padded stride for ssm-conv weights (llama/24470)

ggml : optimize AMX (llama/24806)

Flatten the partition over n_batch * M so every thread participates in
the quantization

    | CPU                             | Model                         | Test   |   t/s OLD |   t/s NEW |   Speedup |
    |:--------------------------------|:------------------------------|:-------|----------:|----------:|----------:|
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B IQ4_NL - 4.5 bpw  | pp512  |    730.71 |    779.86 |      1.07 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B IQ4_NL - 4.5 bpw  | tg128  |     87.88 |     86.79 |      0.99 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B IQ4_XS - 4.25 bpw | pp512  |    725.09 |   1023.31 |      1.41 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B IQ4_XS - 4.25 bpw | tg128  |     83.64 |     83.62 |      1.00 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_0              | pp512  |    820.51 |    924.05 |      1.13 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_0              | tg128  |     90.59 |     92.46 |      1.02 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_1              | pp512  |    776.88 |    872.79 |      1.12 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_1              | tg128  |     89.39 |     90.94 |      1.02 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_K_M            | pp512  |    719.28 |   1009.27 |      1.40 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_K_M            | tg128  |     80.62 |     80.86 |      1.00 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_K_S            | pp512  |    732.29 |   1077.29 |      1.47 |
    | Intel(R) Xeon(R) Platinum 8488C | qwen35 0.8B Q4_K_S            | tg128  |     86.42 |     83.53 |      0.97 |

Signed-off-by: Adrien Gallouët <redacted>

ggml-webgpu: add adapter toggles for F16 on Vulkan + NVIDIA

mtmd, arg: fix utf8 handling on windows (llama/24779)

* mtmd, arg: fix utf8 handling on windows

* also fix ggml_fopen

* fix build fail

* also fix CLI

examples : fix argument flag for min speech duration in VAD (#3907)

Fixed the -vspd flag for vad_min_speech_duration_ms, to prevent hiding vad_min_silence_duration_ms.
In usage () clarified the output timestamp units.
Fixed a few typos.

examples : update model names in parakeet-cli README.md [no ci] (#3906)

This commit updates the parakeet model names which are currently missing
the 'ggml-' prefix.

Fix pkgconfig configuration (Nix build failure) (#3894)

include parakeet in build-xcframework.sh (#3899)

talk-llama : sync llama.cpp

sync : ggml

ggml : bump version to 0.15.2 (ggml/1548)

ggml-cpu: support K tails in power10 Q8/Q4 MMA matmul (llama/24753)

* ggml-cpu: support K tails in Power10 MMA Q8/Q4 matmul

This patch removes the requirement that K be divisible by kc in the tinyBlas_Q0_PPC tiled matmul path. Process the final K panel using its actual depth and pass the reduced panel size through packing and kernel execution.  This allows more workloads to use the MMA kernel and reduces fallback to mnpack.

* Apply suggestion from @taronaeo

Co-authored-by: Aaron Teo <redacted>
---------

Co-authored-by: Aaron Teo <redacted>

Ggml/cuda col2im 1d (llama/24417)

* cuda: add GGML_OP_COL2IM_1D, follow-up to the CPU op

* cuda: col2im_1d use fast_div_modulo for the index decomposition

* cuda: col2im_1d tighten supports_op, type match and contiguous dst

hexagon: support for op-trace (fine-grain tracing of HVX/HMX/DMA events) (llama/24592)

* hex-optrace: add support for optrace and instrument matmul and flash-atten code

* hex-trace: improve trace event and prefetto generator

* hex-trace: add new script dedicated to handling traces, specifically perfetto traces

* hex-trace: add --head/--tail options to profile and trace tools

* hex-trace: fix whitespaces

* hex-trace: fix flake8 warnings

* hex-trace: fix flake8 warnings

* hmx-fa: restore q_tiles clearing

* hex-profile: remove circular dep in includes

* hex-trace: simplify trace sizing check

* hex-profile: sort events in the summary by name

rename GGML_SYCL_SUPPORT_LEVEL_ZERO (llama/24719)

* rename GGML_SYCL_SUPPORT_LEVEL_ZERO to GGML_SYCL_SUPPORT_LEVEL_ZERO_API, and GGML_SYCL_ENABLE_LEVEL_ZERO to  GGML_SYCL_USE_LEVEL_ZERO_API

* fix code format

* fix error when rebase

sycl : support MUL_MAT and OUT_PROD with Q1_0 (llama/24721)

support OPs: conv_2d, conv_2d_dw, conv2d_transpose (llama/24600)

* fix conflict

* fix format issue, rename

* rm debug code

* correct the file name

metal : check for BF16 support in concat kernel (llama/24747)

ggml-cpu: Conditionally enable power11 backend based on compiler support (llama/24687)

* ggml: Conditionally enable power11 backend based on compiler support

Guard POWER11 backend creation behind a compiler flag check for -mcpu=power11. This avoids build failures on current GCC/Clang toolchains while preserving forward compatibility once POWER11 support becomes available.

* Update CMakeLists.txt

ggml-cpu: Use -mcpu=power10 for P10 and P11

metal : implement rope_back operator (llama/24725)

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 (llama/24724)

* 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

add dev2dev memcpy by SYCL API (llama/24476)

* add dev2dev memcpy by SYCL API

* mv GGML_SYCL_DEV2DEV_MEMCPY to runntime table

* update the detect method for p2p comm

* fix the erro created during fix confilct

---------

Co-authored-by: Neo Zhang <NA>

Add conv_3d (llama/24691)

* add conv_3d

* optimize

* update ops.md

* restore test script

* rm unused code

* rm copyright notes

vulkan: record actual memory properties during buffer creation (llama/24326)

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

This reverts commit 0f7fada56bea32c9e0db3874b801d887ff6c4529.

ci: fix vulkan docker images (llama/24595)

* Update vulkan-shaders-gen.cpp

* Update vulkan-shaders-gen.cpp

add comment describing code change intention

* Update vulkan-shaders-gen.cpp

fix potential UB

opencl: optimize mul_mat_f16_f32_l4 for decode (llama/24504)

openvino: OV 2026.2, context-shift, Q5_1 support, gemma4 dense/embedding, and -fa off (llama/24503)

* Add interface is_model_splitted() to check the c-graph is splited or not

* Infer and propagate dynamic-dimension indices for all tensors in the GGML graph in api compute_model_outputs()

* Only do this for fallback sub graph

* Move dynamic dims compute in graph missmatch

* ggml-openvino: fix tensor data handling for PERMUTE/VIEW ops in split models

* ggml-openvino:add comments

* ggml-openvino: override VIEW op_case to 0 for split model inputs

* openvino backend: Handle unsupported VIEW shape-mismatch in OpenVINO backend

* Enable additional mul_mat tests and add tensor data saving function (llama/81)

* ggml-openvino: fix CONT/TRANSPOSE mapping and improve dynamic-dimension handling

* OpenVINO: add NORM/TANH support and rework SOFT_MAX translation

* ggml-openvino: extend VIEW handling

* Enable -fa off (llama/118)

* Enable --context-shift

* Fix llm param compute error for normal softmax not the softmax in attention

* OpenVINO backend: fix error for attention size compute in llm param

* use tensor->extra in infer_request i/o

* OpenVINO backend: refacter the compute_llm_params() func add get_attention_pattern_case to easy extand

* OpenVINO backend: clean unused code

* 1to1 match op update (llama/146)

* added translate_1to1_match_1_input function and updated gelu and tanh translations

* Remove unused translation function calls

---------

Co-authored-by: Mustafa Cavus <redacted>
* initial gemma4 support

* removed hardcoded names for kv cache slicing

* OpenVINO backend: Add new attention pattern for llm parameters compute

* flash attn Q shape static conversion

* Remove slice in permute translation when n_seq is 1

* return optional in extract_layer_from_name

* OpenVINO backend: refactor VIEW related operation (llama/148)

* OpenVINO backend: refactor VIEW related operation

* Enable VIEW handling in following ops

* OpenVINO backend does not support GGML_OP_NORM & GGML_OP_L2_NORM with VIEW input accuracy issue from OpenVINO

* OpenVINO backend: Add ops l2_norm & pad

* OpenVINO backend does not support CPY with non-contiguous data or mismatched types

* add op SSM_CONV GATED_DELTA_NET

* OpenVINO backend: fix error for bf16 in OV gpu plugin

* reverted static Q input shape for attention layer

* OpenVINO backend: remove hardcode name inp_tokens, which ignore some leaf case

* Disable remote tensor due to bug in ov gpu

* Disable n_token > 1 GATED_DELTA_NET on gpu

* OpenVINO backend: fix the view op dynamic handling issue in gemma4 & enable view + get_row

* OpenVINO backend: clean code

* OpenVINO backend: enable view + norm/rms_norm

* OpenVINO backend: concat op

* OpenVINO backend: argsort op

* OpenVINO backend: enable unary + view & GGML_UNARY_OP_SOFTPLUS

* Fix issue for test-backend-ops in TOPK_MOE, which compare VIEW ops result, VIEW node in OpenVINO no need compare, the whole graph result is correct

* OpenVINO backend: enable sum_rows

* OpenVINO backend: enable clamp

* OpenVINO backend: enable DIV

* OpenVINO backend: enable GGML_OP_MUL_MAT_ID

* OpenVINO backend: disable MUL_MAT_ID_FUSION case with large mem needed

* OpenVINO backend: Disable GGML_OP_ARGSORT, cause test_backend-ops failed

* OpenVINO backend: fix issue in mul_mat_id

* OpenVINO backend: Disable DIV with broadcast on GPU

* OpenVINO backend: update DIV

* use ov internal op GatedDeltaNet

* OpenVINO backend: enable llama erch test qwen3next

* OpenVINO backend: enable RMS_NORM + VIEW & remove op_case 2 for rope

* OpenVINO backend: fix error

* suggested changes, need review

* suggested changes, need review

* OpenVINO backend: clean unused code & fix build warning

* OpenVINO backend: enable minicpm3 for arch test

* Disable GDN op (llama/177)

* disable gated_delta_net

* update stateful_kv_size correctly in mismatch case

* OpenVINO backend: enable arch test for qwen3vl

* OpenVINO backend: enable cohere2 for arch test

* OpenVINO backend: enable t5 for arch test

* OpenVINO backend: enable jamba for arch test

* OpenVINO backend: remove warning for tmp

* OpenVINO backend: enable kimi-linear for arch test

* Remove unused

* Fix gpt-oss accuracy issue

* OpenVINO backend: enable arctic for arch test

* OpenVINO backend: enable grok for arch test

* Gemma4 initial npu support (llama/179)

* Initiall gemma4 npu support

* temp. fix for gemma4 accuracy bug on npu

* Remove hardcoded names for npu-fold handling

* revert static n tokens for cont translation as it is not needed

* removed unused variable

* ggml-openvino: add GGML_OPENVINO_ENABLE_CACHE env var to control decoder cache. Add environment variable GGML_OPENVINO_ENABLE_CACHE (default: YES). When set to NO, the decoder_cache is bypassed and models are rebuilt from the cgraph on every inference call in both dynamic and static compute paths. This is useful for debugging and verifying correctness without caching interference.

* Revert "Gemma4 initial npu support (#179)"

This reverts commit 0d29a9c4a52dc2c8aa52990f1a3854cfb01768ad.

* OpenVINO backend: disable debug log print

* Update TBB discovery. Delegated to OpenVINOs own config.

* OpenVINO backend: GGML_OPENVINO_ENABLE_CACHE YES -> 1

* OpenVINO backend: fallback FLASH_ATTN_EXT in gemma3n to CPU backend

* Add raw ov infer profiling metric

* Add OV raw infer time metric to static compute path

Co-authored-by: virajwad <redacted>
* Modify precision of static profiling

* update to OV 2026.2, add OV windows CI

* fix editorconfig-checks

* Initiall gemma4 npu support

* temp. fix for gemma4 accuracy bug on npu

* Remove hardcoded names for npu-fold handling

* revert static n tokens for cont translation as it is not needed

* removed unused variable

* test-llama-archs fix

* Fix gemma4 flash_attn fallback

* support im2col

* fix code style

* disable add_rope_sin_cos optimization

* stateless boradcast and rope optimizations

* Enable manual gqa attn by default for stateless gpu

* manual gqa: fixed static batch

* gemma4 llama-bench ctx update fix

* Update OV win CI

* stateful rope fusion temp. fix

* OpenVINO backend: Conslolidate supported ops

* Exclude unsupported GGML_OP_SUB cases

* Exclude unsupported TOPK_MOE cases

* OpenVINO Backend: MUL_MAT enhancements

* Update OV CI

* support f16 mask input for npu

* Make GGML_OPENVINO_* env vars usage uniform

Standardize all GGML_OPENVINO_* env flags:
positive integers >0 to enable. Unset, empty, =0, or non-numeric values to disable.
This fixes cases where text values or empty strings enabled features.

* OpenVINO backend: Enhance envvar handling

* more cleanup

* move ggml_openvino_env_flag to appropriate place

* OpenVINO backend: add REPEAT translator, Q5_1 weights, and GLU view-input fix

* ggml-openvino: fix -Werror=cast-qual in extract_q5_1_data

* Update openvino.Dockerfile

Use BuildKit cache mounts for faster Docker rebuilds.
Use apt instead of dpkg, remove unused .ddeb downloads, add DLLAMA_BUILD_TESTS=OFF.

* ggml-openvino: centralize env var access via *getenv_str/getenv_int helpers

Replace getenv and legacy flags with _str and _int helpers.Minor cleanup, doc updates.

* OpenVINO backend: Enable GGML_OP_ADD_ID

* Uptade openvino backend clamg-format

* clang-format

* Update OPENVINO.md (llama/211)

* OpenVINO backend: fix accuracy issue for op CONCAT with i64 precision

* Remove strict concurrency for gpu-openvino-low-perf

* Update openvino CI keynames; add ccache-clear

* Apply suggestions from code review

Co-authored-by: Sigbjørn Skjæret <redacted>
* Fix formatting

---------

Co-authored-by: Xuejun Zhai <redacted>
Co-authored-by: Mustafa Cavus <redacted>
Co-authored-by: Mustafa Cavus <redacted>
Co-authored-by: Xuejun <redacted>
Co-authored-by: Wang Yang <redacted>
Co-authored-by: Ravi Panchumarthy <redacted>
Co-authored-by: virajwad <redacted>
Co-authored-by: copilot-swe-agent[bot] <redacted>
Co-authored-by: Mostafa Faheem <redacted>
Co-authored-by: Sigbjørn Skjæret <redacted>

sycl : Enable to support fp16 by OPs: SQR, SQRT, LOG, SIN, COS, CLAMP (llama/24692)

SYCL: fix use-after-free bug with async memcpy in MoE prefill (llama/24676)

* SYCL: fix a bug with async memcpy

* make mmid_row_mapping_host persistent

* comment on stream->wait

* Apply suggestion from @sanmai

* Apply suggestion from @sanmai

* Apply suggestion from @sanmai

sycl: Add optional USM system allocations (llama/22526)

This introduces an optional feature to allocate large GPU buffers (≥ 1GB)
using USM system allocations if supported by the device. It allows using
buffers from the system allocator then letting the system manage memory
migrations between host and device as necessary.

This feature is disabled by default and requires the GGML_SYCL_USM_SYSTEM
environment variable to enable. If USM system allocations are not supported
by the device or the system, we fallback to regular allocations.

This feature can allow VRAM overcommit. For example, the test below fails
on B580 due to lack of memory for allocation, but it passes when enabling
USM system allocations:

  ./examples/sycl/test.sh -m Qwen3.5-27B-Q3_K_M.gguf -lv 4

Signed-off-by: Francois Dugast <redacted>

vulkan: prefer host-visible memory buffers on UMA devices (llama/22930)

* implement UMA host-visible memory

* update based on 0cc4m's suggestion

vulkan: Support gated_delta_net with S_v=16 (llama/24581)

sycl: support reordered Q4_K/Q5_K/Q6_K MoE MUL_MAT_ID (llama/24452)

* sycl: support reordered Q4_K and Q5_K MoE MUL_MAT_ID

Extend reordered-weight handling to fused MoE MUL_MAT_ID for Q4_K and Q5_K expert tensors and add Q5_K reordered DMMV coverage. Unsupported 3D reorder cases now fall back instead of aborting.

* sycl: extend MoE reorder to Q6_K mul_mat_id

Support OP EXPM1, support all UT cases of FLOOR, TRUNC, ROUND (llama/24363)

* support OP EXPM1, support all UT cases of FLOOR, TRUNC, ROUND

* fix conflict

* rebase, support new UT case of repeat, concat

vulkan: add col2im_1d op (llama/24425)

* vulkan: add GGML_OP_COL2IM_1D, follow-up to the CPU op

* vulkan: col2im_1d bounded gather loop instead of full-K scan with modulo

* vulkan: col2im_1d address review from @jeffbolznv

* vulkan: col2im_1d return nullptr for unsupported types, address review from @0cc4m

vulkan: support more CONCAT types (llama/24579)

wasm : fix fallback symbol collision (llama/24639)

SYCL: use native subgroup size for K-quant DMMV (llama/21700)

sycl: fix soft_max_f32 max reduction (llama/24451)

sycl : fix reorder function; add fp32/fp16 in build script (llama/24578)

sycl : enhance set_rows to support q1_0, mxfp4, nvfp4 (llama/24564)

add to support pool_1d, move pool_1d/2d code to pool.cpp/hpp (llama/24584)

* add to support pool_1d, move pool_1d/2d code to pool.cpp/hpp

* update ops.md

Remove per-allocation Level Zero runtime checks (llama/23399)

* [SYCL] Centralize Level Zero detection in ggml_sycl_init

* use the same wording

* get back the warning

* [SYCL] Remove per-allocation getenv() for GGML_SYCL_ENABLE_LEVEL_ZERO

* bring back the comment

* move it up to make sure devices call the shots

* move the env detection early

* replace g_ggml_sycl_enable_level_zero with a direct call to .ext_oneapi_level_zero

* update the comment

* switch back to g_ggml_sycl_enable_level_zero with a sentinel

* remove the check

* Reduce the diff

* reword, move lower

* move things aroudn

* remove forward declaration if favor of a full replace

* pre-cache results of zeDeviceGetProperties

* put ggml_sycl_get_env back

* replace get_sycl_env with ggml_sycl_get_env

* add whitespace back

* Apply suggestion from @sanmai

metal : add repeat bf16 (llama/24638)

CUDA: only support F32/F16 for GGML_OP_REPEAT (llama/24533)

ggml-webgpu: improve i-quants mul_mat performance and speed up prefill (llama/24530)

* Improve prefill speeds for i-quants

* Fix #if defined() usage in preprocessor guards.

vulkan: support non-contig unary/glu ops (llama/24215)

* vulkan: support non-contig unary/glu ops

Change unary/glu ops to pass in all strides and use fastdiv for the index
calculation. Put all unary ops in one file, similar to glu, to share the
code. codex went ahead and added expm1 without me asking, but I had to
make it do a real precision analysis rather than just making stuff up.

unary.comp initially couldn't use generic_unary_head because there wasn't
space for xielu's additional constants. Fixing this required packing the
fastdiv 'L' values.

* attempt to workaround compiler bug

* resolve conflict from #23991

* use expm1

vulkan: add pipeline barriers for memcpy read operations (llama/23770)

* vulkan: add pipeline barriers for memcpy read/write operations

* remove unnecessary host write pipeline barriers

release : v1.9.1 (#3892)

ci : add GGML_NATIVE=OFF and GGML_BMI2=OFF to windows-blas (#3891)

* ci : add GGML_NATIVE=OFF and build all cpu-variants

This commit adds -DGGML_BACKEND_DL=ON, -DGGML_NATIVE=OFF, and
-DGGML_CPU_ALL_VARIANTS=ON to the releases.

The motivation for this is that currently the Windows BLAS build
uses the native CPU instructions and if target systems do not support
these instructions, the build will fail like the linked issue reports.

Resolves: https://github.com/ggml-org/whisper.cpp/issues/3889

* ci : update ubuntu-cpu release job for all variants [no ci]

This commit enables the ubuntu-cpu job to include all cpu variants and
ensures that the ggml backend libraries are built into the bin directory
similar to how llama.cpp does it.

The following is a build on my fork with this change:
https://github.com/danbev/whisper.cpp/releases/tag/untagged-fc3c71f0bf0f7bf19d19

release : v1.9.0 (#3886)

ruby : add support for Parakeet (#3885)

* Add Whisper::Parakeet::Params

* Add tests for Parakeet::Params

* Remove unused variabel

* Add callbacks to Parakeet::Params

* Group callback and user_data params

* Undefine local macros

* Define GetParakeetParams

* Remove unused variable

* Use ITERATE_CALLBACK_PARAMS

* Use ITERATE_CALLBACK_PARAMS instead of ITERATE_USER_DATA_PARAMS

* Fix memsize

* Remove unnecessary macros

* Simplify params registration

* Define Parakeet

* Add hook methods to Parakeet::Params

* Fix typo

* Check callback container in GetParakeetParams

* Reduce if

* Free parakeet_full_params

* Implement Parakeet::Context#initialize

* Add TestParakeetContext

* Add Parakeet::Segment

* Prevent double-free

* Add Parakeet::Context#transcribe

* Add Parakeet::Context#each_segment

* Define Parakeet::Segment attributes

* Define Parakeet::Segment#deconstruct_keys

* Add tests for Parakeet::Segment#deconstruct_keys

* Run Parakeet::Context#transcribe without GVL

* Make it to abort for Parakeet

* Add Parakeet.log_set

* Define Parakeet::Token

* Define Parakeet::Segment#each_token

* Implement some hooks of Parakeet::Params

* Convert int to VALUE

* Implement hooks for Parakeet

* Implement Parakeet::Context#full

* Add tests for Parakeet::Context#full

* Add Parakeet to RBS

* Fix ruby_whisper_parakeet_params_free

* Free ruby_whisper_parakeet_context

* Add tests for hooks

* Add Parakeet section to README

* Add more attributes of Parakeet::Context

* Add tests for Parakeet::Context's attributes

* Update RBS

* Register parakeet-tdt-0.6b-v3

* Narrow scope of log constants

* Extract activate and deactivate of log_queue

* Make start_log_callback_thread private

* Don't call start_log_callback_thread unncecessarilly

* Early return from log_queue_enqueue when not active

* Gropu log_queue members

* is_active -> is_open

* Fix English

* Share parakeet full body function

* ruby_whisper_parakeet_abort_callback_user_data -> ruby_whisper_abort_callback_user_data

* NULL check for callback containers

* Fix Parakeet.log_set

* Omit Parakeet tests on CI

* Extract Whisper::LogSettable

* Join log callback thread in a log queue function

* Revert Join log callback thread in a log queue function

* Extract output methods to modules

* Move Parakeet init functions into init_parakeet()

* Add output methods to Parakeet classes

* Add Parakeet's output methods to RBS

* Use Whisper::Output in RBS

* Add LogSettable to RBS

* Fix module Token -> class Token

* Add Parakeet::Model

* Add test for Parakeet::Model

* Add Parakeet::Model to RBS

* Move position of Parakeet::Model in RBS

* Parakeet -> TestBase::Parakeet

* Add Parakeet::Context#model in RBS

* Add Whisper::Output

* Fix nil check

* Define ruby_whisper_parakeet_model_memsize

* Fix order of declaration in ruby_whisper_parakeet_model_get_xxx

* Define Parakeet.system_info_str

* Add test for Parakeet.system_info_str

* Add signature of Parakeet.system_info_str

* Define Parakeet::VERSION

* Add test for Parakeet::VERSION

* Add signature of Parakeet::VERSION

* Add Parakeet::Context::Params

* Make Parakeet::Context.new accept Context::Params

* Add test for Parakeet::Context.new with Context::Params

* Update RBS

* Remove params from Parakeet::Params which are moved from whisper_parakeet_full_params

* Remove tests for removed params

* Make Parakeet tests follow original behavior changes

* Add Parakeet model shortcuts

* Alloc token data in factory instead of alloc func

* Fix variable name

* Update RBS

* Refactor log settable module

* Use log settable for Whisper

* Address deadlock

* Make test follow change of log queue implementation

* Refactor to make abort callback use the same way to parakeet's way

* Remove redundant structs

* Fix test name

* Fix README

* Add missing parallel transcription

* Fix test for parakeet info

* Remove removed params

* Wait for logs dequeued

* Fix instance variable name

* Load etc feature

* Remove unnecessary comment

* Remove unnecessary thread safety check

* Remove outdated comment

* Skip downloading model if cache exists

* Change Hugging Face URI for Parakeet models

* Bump required Ruby version to 3.3

* Fix English

parakeet : add support for NVIDIA Parakeet (#3735)

* parakeet : add support for NVIDIA Parakeet

Co-authored-by: Georgi Gerganov <redacted>

ci : only trigger release jobs for tags (#3883)

* ci : only trigger release jobs for tags

This commit removes the building of the release jobs on pushed to
master.

The motivation for this is that it can be confusing at the momement when
releasing that the push to master also triggers the release jobs but
the actual release will be skipped. With this change the release job is
only run when a tag is pushed which should result in a single Release
github actions job and make it easier to follow.

* ci : add GGML_NATIVE=OFF for ubuntu-22-gcc

release : v1.8.7 (#3881)

cli : add --version flag (#3878)

Adds a `--version` option to whisper-cli that prints the library version
via `whisper_version()` and exits, plus a corresponding entry in the help
output. Mirrors the existing `-h`/`--help` handling.

Closes #608

talk-llama : sync llama.cpp

sync : ggml

ggml : bump version to 0.15.1 (ggml/1541)

ggml: support concat for scalar types at cuda backend (llama/24011)

* cuda: support concat for scalar types

* Update concat.cu

* fix metal ci issue

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