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 johannesg@5d6.de

• 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 johannesg@5d6.de

• Update ggml/src/ggml-cuda/template-instances/generate_cu_files.py

Co-authored-by: Johannes Gäßler johannesg@5d6.de

Co-authored-by: Johannes Gäßler johannesg@5d6.de

common : re-arm reasoning budget after DONE on new (#22323)

DONE state absorbs all tokens including a new start tag, causing any think blocks after the first to run unbudgeted. Observed on unsloth/Qwen3.6-27B-GGUF which interleaves multiple blocks per response.

Fixed by advancing start_matcher in DONE branch and re-arming to COUNTING with a fresh budget on match. Adds regression test (test-reasoning-budget: test 6).

vulkan: Coalesce Q4_K/Q5_K scale loads (#21751)

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 u8s 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
KhronosGroup/glslang#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

On one Arc Pro B60, unsloth/Qwen3.5-35B-A3B-GGUF:UD-Q4_K_XL:

• pp512: 907.34 ± 9.28 -> 941.94 ± 10.53 (+4%)
• pp2048: 897.95 ± 1.82 -> 931.55 ± 1.79 (+4%)
• tg128: 49.49 ± 0.02 -> 49.86 ± 0.05 (+ <1%)

On one Arc Pro B60, unsloth/Qwen3.5-27B-GGUF:Q4_K_S:

• pp512: 324.13 ± 10.52 -> 354.33 ± 6.81 (+9%)
• pp2048: 329.80 ± 0.25 -> 357.10 ± 0.06 (+8%)
• tg128: 17.11 ± 0.01 -> 18.11 ± 0.01 (+6%)

On four Arc Pro B60s, unsloth/Qwen3.5-122B-A10B-GGUF:Q5_K_S with
-sm layer (note that -sm tensor improvements will naturally be
less):

• pp512: 264.55 ± 2.81 -> 280.45 ± 3.94 (+6%)
• pp2048: 319.32 ± 2.72 -> 335.70 ± 3.48 (+5%)
• tg128: 26.39 ± 0.01 -> 26.67 ± 0.01 (+1%)

ggml-webgpu: fix buffer aliasing for ssm_scan and refactor aliasing logic (#22456)

• Refactor buffer aliasing to be part of shader lib decisions

• cleanup

• formatting

vulkan: add barrier after writetimestamp (#21865)

ggml : skip already registered backends and devices (#22296)

Signed-off-by: Adrien Gallouët angt@huggingface.co

ggml : revert to -lm linking instead of find_library (#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 angt@huggingface.co

• ggml : use MATH_LIBRARY only if defined

Signed-off-by: Adrien Gallouët angt@huggingface.co

• ggml : fix initial broken condition

Signed-off-by: Adrien Gallouët angt@huggingface.co

• ggml : always respect MATH_LIBRARY when defined

Signed-off-by: Adrien Gallouët angt@huggingface.co

Signed-off-by: Adrien Gallouët angt@huggingface.co

CANN: add new ops, optimize existing ops (#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

spec : refactor params (#22397)

• spec : refactor params

• cont : fix

• cont : rename "sparam" to "sampling"

• cont : add spec params category

• cont : add info about removed arguments

• cont : skip param length check for spec params

• cont : adapt server tests

server: use pos_next instead of n_tokens for m-rope (#22439)