shalinib-ibm [Fri, 19 Jun 2026 05:55:38 +0000 (11:25 +0530)]
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.
shalinib-ibm [Wed, 17 Jun 2026 18:45:19 +0000 (00:15 +0530)]
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.
Georgi Gerganov [Wed, 17 Jun 2026 16:38:55 +0000 (19:38 +0300)]
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
* 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.
* 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
Francois Dugast [Wed, 17 Jun 2026 05:54:21 +0000 (07:54 +0200)]
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:
Frosty40 [Tue, 16 Jun 2026 05:35:00 +0000 (00:35 -0500)]
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.
Jeff Bolz [Sat, 13 Jun 2026 13:44:15 +0000 (08:44 -0500)]
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.
Daniel Bevenius [Thu, 18 Jun 2026 12:49:08 +0000 (14:49 +0200)]
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.
* 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
Daniel Bevenius [Tue, 16 Jun 2026 12:33:42 +0000 (14:33 +0200)]
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.
Rum Nguyen [Tue, 16 Jun 2026 06:58:09 +0000 (13:58 +0700)]
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.
Gaurav Garg [Wed, 10 Jun 2026 17:51:16 +0000 (23:21 +0530)]
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.
Oliver Simons [Wed, 10 Jun 2026 12:27:08 +0000 (14:27 +0200)]
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
Pascal [Tue, 9 Jun 2026 09:01:37 +0000 (11:01 +0200)]
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.
Jeff Bolz [Mon, 8 Jun 2026 08:40:37 +0000 (03:40 -0500)]
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.
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.
Mason Milburn [Fri, 5 Jun 2026 05:10:31 +0000 (01:10 -0400)]
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.
Kartik Sirohi [Thu, 4 Jun 2026 13:12:38 +0000 (18:42 +0530)]
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
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
Andreas Kieslinger [Wed, 3 Jun 2026 11:56:42 +0000 (13:56 +0200)]
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>
---------
Max Krasnyansky [Tue, 2 Jun 2026 06:40:08 +0000 (23:40 -0700)]
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
overview / pkg / ggml / sources / whisper.cpp / log