Devedse [Sat, 9 May 2026 05:50:24 +0000 (07:50 +0200)]
Add BF16 support to GET_ROWS operation (llama/21391)
Add GGML_TYPE_BF16 to the SYCL backend's GET_ROWS operation, both in
supports_op and in the kernel dispatch. This fixes a performance
regression where models using BF16 embedding tensors (e.g., Gemma4's
per_layer_token_embd.weight) fall back to CPU for the GET_ROWS op,
causing a full GPU-to-CPU tensor transfer every token.
The fix reuses the existing get_rows_sycl_float template with
sycl::ext::oneapi::bfloat16, matching the pattern already used for
sycl::half (F16) and float (F32).
Signed-off-by: Chun Tao <redacted>
* Remove unneeded/unnecessary comments and annotations
The MMQ subgroup annotations added are on functions gated behind
ggml_sycl_supports_mmq(). Revisit the need for these annotations
when that function changes.
---------
Signed-off-by: Chun Tao <redacted> Co-authored-by: Chun Tao <redacted> Co-authored-by: Todd Malsbary <redacted>
AesSedai [Sat, 9 May 2026 03:28:29 +0000 (20:28 -0700)]
Add flash attention MMA / Tiles to support MiMo-V2.5 (llama/22812)
* mimo-v2.5: add flash attention mma/tiles for for d_kq=192 d_v=128
* mimo-v2.5: follow (256, 256) fattn templates
* mimo-v2.5: cleanup comments
* mimo-v2.5: further comment cleanup
* mimo-v2.5: address PR feedback
fix GQA handling
check for other dangling 320/576 carveouts and mirror them for 192
Add to backend ops test so new paths are covered
Yanzhao Wang [Sat, 9 May 2026 00:12:04 +0000 (17:12 -0700)]
hexagon: add HTP kernel for GGML_OP_GATED_DELTA_NET (llama/22837)
Implement the Gated Delta Net recurrence on HVX with:
- 4-row fused kernels for PP (prompt processing) path
- 8-row fused kernels for TG (token generation) path, reducing
K/Q/gate vector reload overhead by 2x
- Separate PP/TG thread functions for I-cache isolation
- VTCM state scratchpad with DMA in/out for TG single-cycle access
- Vectorized gate exp via hvx_exp_f32
Pascal [Fri, 8 May 2026 09:44:09 +0000 (11:44 +0200)]
cuda: fuse snake activation (mul, sin, sqr, mul, add) (llama/22667)
* cuda: fuse snake activation (mul, sin, sqr, mul, add)
Add ggml_cuda_op_snake_fused with F32 / F16 / BF16 templates. The
matcher recognizes the naive 5 op decomposition emitted by audio
decoders (BigVGAN, Vocos) for snake activation
y = x + sin(a*x)^2 * inv_b and rewrites it to a single elementwise
kernel.
Add test_snake_fuse comparing CPU naive vs CUDA fused across
F32 / F16 / BF16.
* cuda: address review feedback from @am17an
Use ggml_cuda_cast for F32/F16/BF16 conversions and rename
kernel_snake to snake_kernel to match upstream conventions.
* cuda: snake fusion fastdiv on T_len, Suggested-by: @am17an
Signed-off-by: Chun Tao <redacted> Signed-off-by: Todd Malsbary <redacted> Co-authored-by: Chun Tao <redacted> Co-authored-by: Todd Malsbary <redacted>
Bjarke Viksøe [Sun, 10 May 2026 14:24:12 +0000 (16:24 +0200)]
whisper : fix incorrect timestamps, usually near silences (#2279)
* Incorrect timetstamps
Fixes #2271
- Adds consecutive timestamps after end of last segment as the new starting ts
- Add these timestamp to output when "print-special" enabled
- Fixes fflush usage in live reporting
I was not able to test this with the special "token_timestamps" option.
Yiwei Shao [Sat, 2 May 2026 03:29:13 +0000 (20:29 -0700)]
hexagon: hmx flash attention (llama/22347)
* hmx: extract shared interleave headers and unify matmul batched
* hmx: add HMX-accelerated flash attention for prefill
* hmx: replace asm wrappers with Q6_ intrinsics in hmx-utils.h
Switches three single-instruction helpers from inline asm to the matching
Q6_ intrinsics, matching the style established by aizip f8737609a and used
by the upstream PR #21554 hmx-matmul-ops.c rewrite:
hmx_load_tiles_fp16 stays on inline asm: it uses ":deep" activation
streaming, and the mixed Q6_activation_hf_mxmem_RR_deep + non-deep
Q6_weight_hf_mxmem_RR pair fails the HMX backend constraint check
("activate weight pair (1) exceeds limit (1)"). The asm bundle keeps
both halves in one VLIW packet and avoids the diagnostic.
Functionally equivalent — same instructions emitted; the Q6_ intrinsics
just give the compiler more visibility for scheduling.
* hmx: drop the duplicate interleave_fp16_weight_chunk_to_tiles
* hmx: apply upstream optimization to hmx-flash-attn-ops.c
apply restrict, __builtin_assume, and pointer accumulation to the three HMX workers (qk_dot, o_update, o_norm) and the matching inline HMX loops in op_hmx_flash_attn_ext.
* hmx: unify interleave helper
* hmx: multi-thread Q load / O store and enable prefill FA dispatch
Extract inline Q-load and O-store loops into worker_pool-parallel helpers
(fa_phase_q_load, fa_phase_o_store) so HVX threads split the F32↔F16
conversion work across row ranges. Also relax the softmax threading
gate from n_row_vec_cnt >= n_threads to >= 2, which was unnecessarily
forcing single-thread fallback when n_rows_g < 512.
On the dispatch side, remove the ne[2] != 1 guard that blocked multi-head
(prefill) FA from reaching the HTP backend — GQA is already handled
internally by both the HMX and HVX flash-attention paths.
* hmx: relax matmul pipeline gate to cover k > n shapes (e.g. FFN_down)
* hmx: optimize FA softmax mask phase (no-ALiBi fast path + GQA dedup)
* hmx: Add an asm memory clobber at the phase boundary to prevent reorder bug
* [experimental]: fp16 softmax (EXP2_HF) to accelerate fa
Bake log2(e) into qk_scale and use hvx_exp2_hf directly for P and m_diff
(base-2 consistent, matches htp-ops-lib). ~22 ALU ops for 64 lanes vs
~44 for the F32 round-trip path.
* hmx flash-attn: refine cost model coefficients based on profiling data
* hmx flash-attn: replace asm clobber with targeted volatile reads on vtcm_d_tiles
* hmx flash-attn: preserve additive mask bias in no-ALiBi fast path
The no-ALiBi fast path (max_bias==0) was skipping mask add entirely on
the assumption that mask values are only {0, -inf}. This is wrong when
the mask carries additive positional bias — those terms were silently
dropped. Keep the slope-mul skip (slope≡1.0) but add mask back so the
bias survives; vmux still clamps below -16 to -inf.
Also add HMX FA coverage to test-backend-ops: prefill shapes (nb=64,
nb=32) × {mask on/off} × {ALiBi on/off} × {softcap on/off}, F16 KV,
hs ∈ {64, 128}.
- flash-attn: when EXP2_HF is on AND logit_softcap is active, fold
log2(e) into the post-tanh multiplier (v_cap) instead of pre-baking
it into qk_scale. Pre-baking shifted the tanh knee from x≈c to
x≈c/log2(e) and produced numerically wrong softcapped outputs
whenever both knobs were enabled.
- flash-attn softmax (fa_softmax_thread): replace the union+memcpy
scalar extract pattern with HVX vmux-based per-row accumulators on
rowmax/rowsum. Add hvx_vec_get_f16 helper in hvx-base.h. Functional
parity, less scalar code, clearer hf/qf16 lane-format contract.
- matmul (hmx_mat_mul_permuted_qk_0_d16a32): pick pipeline vs sequential
layout based on whether the chunker actually yields >=2 n-chunks,
instead of the static (m>=128 && n>=256) gate. Avoids paying for
output double-buffer + worker dispatch when there is no HMX/HVX
overlap to gain (e.g. shapes that collapse to one n-chunk).
- tests: add HMX flash-attention coverage over the
{mask, ALiBi (max_bias), logit_softcap} cross-product for the prefill
path — head_dim 64/128, GQA 4×4, kv=512/nb=64 plus a kv=113/nb=32
non-aligned case.
* [Help Wanted]: refactor D matrix computation into separate function for clarity and maintainability
* format code
* hexagon: looks like -O3 is causing issues with the large code base, switch to -O2 and -flto instead
* hexagon: use hex_ prefix for swap_ptr
* hexagon: move vtcm_seq_alloc into vtcm-utils.h
More vtcm allocator updates are coming so it makes sense to start the separate hdr for it.
* hmx-utils: add hmx_prefix for layout converters
* hmx-mm: move main hmx_mm functions to the end, remove unused fwd decls, etc
* hmx-mm: remove unused qweight_fetch_task_state_t and minor alignment fixes
* hmx-fa: minor alignment fixes
* hmx-fa: move hmx_flash_atten into hmx-ops.h
* hmx-fa: remove redundant workpool pointer in the hmx_fa_ctx, plus minor alignment updates
* hmx-fa: minor alignment and simplifications
* hexagon: move FA_EXP_F16 option to hostside CMake file
* hmx-fa: use hvx_vec_splat_f16 instead of fp16_to_bits
* hmx-fa: add hvx_splat_u16/u8 and use that in the fa instead custom hvx_fill
* hmx-fa: some more alignment updates in the core fa function
* hmx-fa: keep slopes in vtcm in fp16
Saves malloc/free and removes the need for float -> fp16 downcast on every use.
Jeff Bolz [Fri, 1 May 2026 13:28:32 +0000 (15:28 +0200)]
vulkan: Support asymmetric FA in coopmat2 path (llama/21753)
* vulkan: Support asymmetric FA in coopmat2 path
There has been some recent interest/experimentation with mixed quantization
types for FA. I had originally designed the cm2 FA shader with this in mind
(because I didn't realize it wasn't supported at the time!), this change
adds the missing pieces and enables it.
Also support Q1_0 since people have been trying that out (seems crazy, but
who knows).
We should be able to do similar things in the coopmat1/scalar path, but
there's another change open against the scalar path and I don't want to
conflict.
ggml-cpu : disable tiled matmul on AIX to fix page boundary segfault (llama/22293)
* ggml-cpu : disable tiled matmul on AIX to fix page boundary segfault
vec_xst operations in the tiled path crash on AIX when writing
near 4KB page boundaries due to strict memory protection. Fall
back to mnpack implementation on AIX for stable execution.
ggml-cpu: cmake: append xsmtvdotii march for SpacemiT IME (llama/22317)
* ggml-cpu: cmake: append xsmtvdotii march for SpacemiT IME
When GGML_CPU_RISCV64_SPACEMIT=ON is set, ime1_kernels.cpp contains
inline asm for the vmadot family which requires the xsmtvdotii custom
extension.(problem can see in some blogs and make sure in K3 platform)
The current CMakeLists does not include xsmtvdotii, so any toolchain
that honours the explicit -march (tested with SpacemiT GCC 15.2) fails
at the assembler stage:
Append _xsmtvdotii to MARCH_STR when GGML_CPU_RISCV64_SPACEMIT is
enabled so the IME path can actually build with a capable toolchain.
No effect on builds that leave GGML_CPU_RISCV64_SPACEMIT off.
toolchain from https://www.spacemit.com/community/resources-download/Tools
ggml-cuda: add flash-attn support for DKQ=320/DV=256 with ncols2=32 (… (#22286)
* ggml-cuda: add flash-attn support for DKQ=320/DV=256 with ncols2=32 (GQA=32)
Adds MMA-f16 and tile kernel configs, dispatch logic, template instances,
and tile .cu file for Mistral Small 4 (head sizes 320/256), restricting to
ncols2=32 to support GQA ratio 32 only.
* Adding check to return BEST_FATTN_KERNEL_NONE in case GQA!=32
* Apply suggestions from code review
Address review comments
Co-authored-by: Johannes Gäßler <redacted>
* Address review comments and making kernel config default to DQK=512, DV=512 instead of DQK=256,DV=256
* Fixed bug with sinks=1, with ncols=32, there are two warp-groups created but sinks index is same(0,...,15) for both the groups hence with sinks=1, output is not matching with CPU output. Added sink_base which will be base index for each warp_group (threadIdx.y / np)
* Apply suggestions from code review
Co-authored-by: Johannes Gäßler <redacted>
* Update ggml/src/ggml-cuda/template-instances/generate_cu_files.py
Co-authored-by: Johannes Gäßler <redacted>
---------
Some SPIR-V compilers (notably mesa) don't handle the current
vulkan Q4_K/Q5_K scale load pattern in mul_mat particularly well.
While reading three `u8`s from the 12-byte scale array should (at
least on some hardware) result in loading the full 12 bytes in a
single LOAD followed by whatever extraction is needed, at least
the ANV Intel driver really can't practically perform this
optimization.
`mesa`'s unsigned upper bound logic doesn't handle tracking bounds
through ternary, resulting in the `(is < 4) ? ... : is - 4` having
an infinite upper bound (as it cannot prove `is - 4` doesn't
underflow). While this could still be rectified if mesa looked at
the array bounds, it currently doesn't and `glslc` currently emits
SPIR-V that doesn't allow for this optimization anyway (though
maybe it will at some point, see
https://github.com/KhronosGroup/glslang/issues/4206).
In mul_mat_vecq we took a different approach to loading the same
fields. We read the first two bytes we needed from `scale` then
took a branch before deciding whether we needed to read a third
byte. In mesa this did, indeed, lead to a top-level branch with
conditional loads. As such these loads ended up not being
coalesced either (at least in the ANV driver) resulting in
additional instructions in our hot loop.
Instead, here, we go ahead and force loading the full 12 bytes and
extract the bits we need from the packed-u32s instead. In mul_mat
there's a few less ternaries and only one extra shift, so even on
drivers that did optimize the previous loads properly the only
material change should be pulling a few extra bytes into registers
(which on most hardware won't cost anything anyway, though
ironically on Intel it theoretically could). In mul_mat_vecq this
requires a bit of extra math and may read bytes from the u32 that
weren't needed, but it seems likely avoiding the branch is a win
on most platforms.
On Intel Xe2/mesa 26.0.4 with the optimizations from
https://gitlab.freedesktop.org/mesa/mesa/-/work_items/15162,
for shader matmul_id_subgroup_q4_k_f32_f16acc_aligned_l:
* Instruction Count: 2753 -> 2688
* SEND Count: 269 -> 261
* Cycle Count: 273976 -> 266138
* Max live registers: 248 -> 246
* Non SSA regs after NIR: 381 -> 382
for shader matmul_id_subgroup_q5_k_f32_f16acc_aligned_l:
* Instruction Count: 2767 -> 2702
* SEND Count: 271 -> 263
* Cycle Count: 274140 -> 268144
* Max live registers: 248 -> 246
* Non SSA regs after NIR: 381 -> 382
for shader mul_mat_vec_id_q4_k_q8_1_f32:
* Instruction Count: 1930 -> 1646
* SEND Count: 116 -> 71
* Cycle Count: 1348306 -> 843350
* Max live registers: 78 -> 84
* Non SSA regs after NIR: 300 -> 135
for shader mul_mat_vec_id_q5_k_q8_1_f32:
* Instruction Count: 2207 -> 1922
* SEND Count: 131 -> 86
* Cycle Count: 1392012 -> 1037836
* Max live registers: 90 -> 90
* Non SSA regs after NIR: 300 -> 135
for shader mul_mat_vec_q4_k_q8_1_f32:
* Instruction Count: 2029 -> 1749
* SEND Count: 111 -> 66
* Cycle Count: 1347278 -> 840118
* Max live registers: 74 -> 80
* Non SSA regs after NIR: 299 -> 134
for shader mul_mat_vec_q5_k_q8_1_f32:
* Instruction Count: 2307 -> 2022
* SEND Count: 126 -> 81
* Cycle Count: 1379820 -> 954042
* Max live registers: 86 -> 86
* Non SSA regs after NIR: 299 -> 134
ggml : revert to -lm linking instead of find_library (llama/22355)
* ggml : revert to -lm linking instead of find_library
`find_library(MATH_LIBRARY m)` was introduced recently, but it breaks
CUDA compilation with GGML_STATIC. I could not find any valid use case
where we would prefer `find_library` over the standard `-lm` approach.
This commit is also meant to start a discussion if there is a valid
reason to keep `find_library(MATH_LIBRARY m)`, we should clarify what
problem it was solving and find an alternative fix that does not break
CUDA with GGML_STATIC.
Signed-off-by: Adrien Gallouët <redacted>
* ggml : use MATH_LIBRARY only if defined
CANN: add new ops, optimize existing ops (llama/21204)
New operators:
- GGML_OP_SET: implement via aclnnInplaceCopy on target region
- GGML_OP_CUMSUM: implement via aclnnCumsum
- GGML_OP_FILL: implement via aclnnInplaceFillScalar
- GGML_OP_DIAG: implement via aclnnInplaceCopy on diagonal strides
- GGML_OP_TRI (lower/lower_diag/upper_diag/upper): implement via
aclnnTril(-1/0) and aclnnTriu(0/1) with appropriate diagonal offsets
- GGML_OP_SOLVE_TRI: implement via aclnnTriangularSolve
- GGML_UNARY_OP_SOFTPLUS: implement via aclnnSoftplus
Optimizations:
- GLU (SwiGLU/GeGLU/GeGLU_ERF/GeGLU_QUICK): fuse with aclnnSwiGlu /
aclnnGeGluV3 when applicable; fallback conditions now checked inside
each function rather than at the call site
- CROSS_ENTROPY_LOSS: replace 5-kernel sequence (LogSoftmax→Mul→
ReduceSum×2→Muls) with single aclnnSoftmaxCrossEntropyWithLogits call
- L2_NORM: fix in-place ClampMin on norm result (was clamping wrong
tensor); add eps clamping before division to avoid divide-by-zero
- PAD_REFLECT_1D: eliminate per-ne[3] loop; assert contiguity and call
ReflectionPad1d once on the full 4-D view; remove redundant nb copies
- GET_ROWS: replace IndexSelect with GatherV2 per batch slice; refactor
helper into gather_batched lambda with batch loop inlined
- SET_ROWS: replace IndexCopy with InplaceIndexCopy per batch slice;
refactor helper into scatter_batched lambda with batch loop inlined
- OUT_PROD: replace O(ne[3]*ne[2]*ne[1]) Ger+InplaceAdd loop with
per-slice Matmul loop (src0 @ src1^T); handles strided-broadcast
batch dims where ne02/ne03 may differ from ne2/ne3
- backend memset_tensor: implement via aclrtMemset (was NULL)
Bug fixes:
- COUNT_EQUAL: use non-inplace EqTensor into a same-type temporary
buffer instead of InplaceEqTensor, avoiding corruption of src0
- ACL graph cache (USE_ACL_GRAPH): restore node_type and src_type[]
fields in ggml_graph_node_properties; has_matching_properties() was
missing type checks, causing F16 and BF16 tensors (same nb[0]=2) to
incorrectly share cached graphs and produce wrong results (ERR≈679)
- graph cache op_params matching: compare full GGML_MAX_OP_PARAMS
bytes so that ops differing only in parameters are not incorrectly
replayed from cache
Fix recurrent state serialization for partial reads and writes (llama/22362)
The previous code worked only for full tensor reads and writes and was hitting `GGML_ASSERT(size == ggml_nbytes(tensor)); ` assert when tested with llama-server.
metal : optimize Metal Tensor API usage for GGML_OP_MUL_MAT (llama/20962)
* Optimize Metal Tensor API usage for matmul2d
Separates the Metal Tensor API (matmul2d) path in kernel_mul_mm into its own standalone kernel, gated by GGML_METAL_HAS_TENSOR.
The legacy simdgroup_matrix kernel is preserved under #else.
Previously both paths were interleaved via #ifdef blocks within a single kernel, forcing the tensor path to share the legacy kernel's data layout and threadgroup memory scheme. Splitting the kernel enabled memory and dispatch optimizations that weren't possible when the two paths shared code structure.
overview / pkg / ggml / sources / whisper.cpp / log