Overview

ggml v0.29.0 adds an OCP FP8 E4M3 tensor type, a new ggml_flash_attn_ext_rows op where each query slice attends to its own list of K/V rows, a batched tensor copy API, and initial support for multi-node tensor parallelism through a new multi-process communicator interface. Core improvements include bounding tensor views by the bytes their strides reach, zero-copy GGUF export of host-buffer tensors, and vm_allocate-based allocations on Apple platforms. Performance work spans all major backends: new SVE kernels in the CPU backend, DPAS flash attention for quantized KV in SYCL, fused MoE MMQ in CUDA, faster MXFP4 mat-vec in Metal, split-K tuning in Vulkan, and a large expansion of quantized GEMM/GEMV coverage in OpenCL. The Hexagon backend gains support for several new ops, including the lightning indexer, top-k mask, conv2d and upscale.

API changes

• Add GGML_TYPE_F8_E4M3 (OCP FP8 E4M3) with float conversion helpers and GGML_FTYPE_MOSTLY_F8_E4M3 (#27512)
• Add ggml_flash_attn_ext_rows(): FA variant where each q slice attends to its own K/V row list via index tensor (#29510)
• Add batched ggml_backend_tensor_copy_n(), ggml_backend_tensor_set_n() and ggml_backend_tensor_get_n() APIs (#30881)
• Add a multi-process communicator interface (ggml_backend_comm_init_rank_t) for collectives across separate ranks (#28967)

Core changes

• Bound tensor views by the bytes their strides reach and validate strided views against the source extent (#31584)
• GGUF writer: write host-buffer tensor data directly to the output stream without an intermediate copy (#31144)
• Use vm_allocate instead of posix_memalign for large allocations on all Apple platforms (#31437)
• Meta backend: keep tensor splits across permuted views and avoid mirrored splits for ARANGE (#30499 #30878)
• Speed up IQ4 quantization with pre-computed halfway points (#31201)

Backend changes

CPU

• Add SVE 256-bit kernels: MXFP4 GEMV (#28141), Q6_K GEMM (#27472) and q8_0 4x8 GEMV (#30585)
• Gate RVV code on the target vector extension so it is not built for rv64gc without V (#31607)
• Implement ggml_flash_attn_ext_rows (#29510)

CUDA

• Add NCCL support for multi-node tensor parallelism via the multi-process communicator, used by the RPC server (#28967)
• Gather batched tensor reads into mapped host memory for the batched copy APIs (#31391)
• Fuse MoE (MUL_MAT_ID) gate/up projections with GLU in MMQ (#31188)
• Implement ggml_flash_attn_ext_rows (#29510)
• HIP: enable sparse attention on CDNA (#31259) and add RDNA4_M family for gfx1170/gfx1171 (#31373)
• Fix Q8_1 correction sum in MMQ (#31247) and premature eviction of CUDA graphs (#30662)

SYCL

• Add DPAS flash attention kernel for q4_0/q8_0 KV caches on XMX (#31125)
• Vectorize reordered-layout dequantization stores for Q4_0, K-quants, IQ3 and MXFP4 (#30976)
• Add fast ggml_sycl_comm_exchange() for tensor parallelism (#30956)
• Add work-group argmax for short TOP_K rows (#31110)
• Fix f32 -> q2_0/q4_0/q5_0 rounding in CPY (#31120) and FLASH_ATTN_EXT failures (#30761)

Metal

• Faster MXFP4 mat-vec; launch unary/binary/GLU kernels with full threadgroups (#31075)
• Implement ggml_flash_attn_ext_rows (#29510)
• Fix cpy kernel for destinations non-contiguous in dim 0 (#31581)

Vulkan

• Add tensor-parallel allreduce from shared host buffers (#25051)
• Improve small-n split-K selection on RDNA3/RDNA4 (#30146) and avoid redundant scans when zeroing skipped rows (#31266)
• Skip the maxStorageBufferRange pre-check for non-compute tensors (#31114)
• Fix out-of-bounds NVFP4 UE4M3 scale lookup (#31079)

OpenCL

• Add quantized GEMV/GEMM support for Q2_0, IQ1, IQ2, IQ3 and IQ4_XS (#31340 #31341 #31267 #31338 #31339)
• Add tiled Q4_K GEMM and faster Q2_0 GEMV for Adreno (#30412 #31577)
• Add FWHT-accelerated GEMM (#31434) and TOP_K_MASK support (#31441)

Hexagon

• Add support for the LIGHTNING_INDEXER (#31289), TOP_K_MASK (#31552), UPSCALE (#31331), CONV_2D (#30995) and DSV4_HC_PRE/POST (#30953) ops
• Faster TOP_K sampler (#31279), block-based HVX IM2COL (#31050), tiled HVX permute (#31446) and better DDR use in transposed CONCAT (#31277)
• Fix DMA transfers with strides larger than 24 bits (#31442)

Other backends

• ZDNN: share buffers with host (#31481) and fix a fattn test failure (#31583)
• WebGPU: add ARANGE support (#31455)
• OpenVINO: slice view inputs of SCALE (#31334)

More info

• Releases and versioning of ggml-org projects

Changelog since v0.28.0

57e2c2c ggml : bump version to 0.29.0 (#1680)
68d509f scripts : add ci-status.sh
689e0cb sync : llama.cpp
f11cb47 ggml : Gate RVV code on the target extension (llama/31607)
5c9c721 sync : llama.cpp
71b8415 ggml : fix bf16 (llama/31609)
b8f6efc opencl: improve q2_0 GEMV for Adreno (llama/31577)
682dfdd sync : llama.cpp
8bdd9c3 ggml: bound views by the bytes their strides reach (llama/31584)
02dcc21 ggml : write host-buffer tensor data directly in gguf writer (llama/31144)
0c4fecd hexagon: add LIGHTNING_INDEXER support (llama/31289)
492ed92 hexagon: add TOP_K_MASK op (llama/31552)
d88c3d4 hexagon: faster TOP_K for the backend sampler (llama/31279)
c8e848e hexagon: add DSV4_HC_PRE and DSV4_HC_POST ops (llama/30953)
6496f89 ggml-zdnn: fix fattn test failure (llama/31583)
781e018 ggml: avoid signed overflow in ggml_roll bounds check (llama/30375)
5d3c669 cuda: potential fix for Q8_1 correction sum (llama/31247)
a044cd3 cuda : gather batched tensor reads into mapped host memory (PR to master) (llama/31391)
40c2633 Add nccl support for multi-node tensor parallelism (llama/28967)
1030236 meta: support GLM-5.3 next sm tensor (llama/30499)
7117df3 metal : fix cpy kernel for dst non-contiguous in dim 0 (llama/31581)
b4134b7 vulkan: skip maxStorageBufferRange pre-check for non-compute tensors (llama/31114)
f9be71f delta_net_base: refactor conv state (llama/30878)
c44f602 vulkan: fix out-of-bounds NVFP4 UE4M3 scale lookup (llama/31079)
e136ad2 cuda : prevent premature eviction of graphs about to be used (llama/30662)
43e15d2 opencl: add initial iq4_xs support (llama/31339)
cb27aa1 webgpu: add ARANGE support (llama/31455)
e96dee8 hexagon: tiled hvx permute (llama/31446)
c33c20d llama : flash_attn_ext_rows - more efficient unified kv for multi-seq (llama/29510)
842f76b ggml-zdnn: share buffers with host (llama/31481)
643b9af ggml : use vm_allocate instead of posix_memalign on all Apple platforms (llama/31437)
2956e59 opencl: add topk_mask (llama/31441)
aa632fc vulkan: add tensor-parallel allreduce function (llama/25051)
122edda vulkan: improve small-n spli...

Overview

This release adds a fused sparse-attention masking operation (GGML_OP_TOP_K_MASK) with CPU, CUDA and Metal implementations, plus accessors to read the precision configured on a compute node. MoE and sparse-attention dispatch got cleaner semantics: an index of -1 in MUL_MAT_ID / ADD_ID now means "skip this slot", so backends skip the work and zero the output row instead of relying on out-of-range index patching. The core relaxes SCALE to non-contiguous views and SUM to non-contiguous inputs. Backends gained a lot of throughput: MoE top-k and gate/up+GLU fusions and chunked graph launches on CUDA, graph record/replay and XMX matrix-engine kernels on SYCL, cooperative GEMMs, Q2_K/Q3_K weights and a lightning indexer on OpenCL, subgroup reductions and src1 precision control on Vulkan, and a restructured HTP backend with new repacking and fusion machinery on Hexagon. ZDNN gained flash attention and WebGPU gained Q2_0.

API changes

• New GGML_OP_TOP_K_MASK op and ggml_top_k_mask() - fused top-k sparse attention mask over picks, blocks and extra cells llama/30640
• New ggml_prec_get_acc() / ggml_prec_get_src() to query the accumulator and per-source precision of a compute node llama/30656
• ggml_mul_mat_id and ggml_add_id accept -1 as an index: the slot is skipped and the output row is zeroed llama/26631
• RPC protocol patch version bumped to 8.0.1 for the new op (GGML_OP_COUNT is now 102)

Core changes

• ggml_scale now accepts non-contiguous row views instead of requiring contiguous input llama/31146
• SUM supports non-contiguous inputs on CPU, CUDA, Vulkan and WebGPU llama/28851
• MUL_MAT_ID / ADD_ID skip the work for -1 indices in all major backends llama/26631
• Meta backend: per-expert token scaling for MoE down projections, plus split-state support for the new masking op llama/30861 llama/31081
• Fix potential counter overflow in ggml_time_ms / ggml_time_us on Windows llama/30969
• ARGMAX added to the ops that may expand during allocation, avoiding undersized buffers llama/30862
• RPC: handle oversized RDMA transfers and std::length_error on message receive instead of aborting llama/30399 llama/31204

Backend changes

CPU

• New CONCAT implementation as a single row kernel covering all types llama/31179
• KleidiAI: Q8_0 SVE dotprod GEMV and SVE i8mm GEMM kernels, with runtime SVE-i8MM feature detection llama/30808
• Vectorized v_silu and v_expf, NEON fp32->fp16 conversion, fp32 gelu_quick/geglu_quick without the fp16 table llama/30791 llama/30505 llama/30737
• Reduce AVX2 tiled accumulators together to cut register pressure in the tiled matmul path llama/30296

CUDA / HIP

• MoE: fused grouped-experts top-k, plus a dedicated fusion path for large expert counts (up to 384) llama/29181 llama/31288
• Fuse the FFN gate/up matmuls with GLU in MMQ using interleaved gate/up rows; fuse MRoPE with SET_ROWS llama/29948 llama/31025
• Chunk CUDA graph launches to amortize the launch overhead; clear graphs after a legacy memory pool flush llama/30804 llama/26771
• Faster MMVQ for low-bit types; MMQ Stream-K capped at available work; DP4A MMQ tile height aligned to the warp size llama/30517 llama/30553 llama/30837
• SSM_CONV supports kernel size 7; RMS_NORM variants loading float4 and keeping x in registers llama/30873 llama/29720
• HIP: byte-permute Q1_0 unpacking with a portable fallback; fix bf16 WMMA operand type for gfx11/gfx12 llama/29927 llama/30915 llama/30885
• Fix shared-memory WAR hazard in vectorized flash attention and a tile-width selection bug in the matmul reduction llama/30573 llama/30836
• Limit the maximum pinned host buffer size to 4 GB llama/30449

Metal

• Implement TOP_K_MASK llama/31230
• Optimized ARGMAX and POOL_1D kernels; removed the amax rescaling from the MoE matmul llama/30862 llama/30924
• Fix 32-bit overflow in large matrix output offsets llama/30072
• Fix missing MoE kernel variant for 9 experts per token, and a moe_reduce fusion race when dst aliases its inputs llama/30699 llama/30795
• Fix memory leaks and cleanup logic in the GPU capture path llama/30890

Vulkan

• Support the src1 precision attribute on matmul ops, configurable through the standard precision setters llama/30656
• RMS_NORM rewritten with subgroup reductions llama/29882
• Small matmul tiles now use one warp tile per subgroup; larger conv3d tile for coopmat2; clamp coopmat1 int8 prefetch llama/30106 llama/29920 llama/30283
• Allow dim-3 broadcast in MUL_MAT and fix gpu-visible host memory handling llama/31239 llama/30468

SYCL

• Add graph record and replay to reduce per-op launch overhead llama/28725
• Restore ESIMD kernels on the FFN path and rework multi-column matmul on Intel XMX matrix engines llama/30239 llama/29864
• Faster dequantization: fast Q8 path, wide stores for Q4_K/Q5_K, merged fp16/fp32 dequant switches, templated MMVQ/concat launches llama/30775 llama/29696 llama/30740 llama/31013 llama/31008
• set_tensor_2d for faster tensor-split (KV-cache) loading; bound weight-reorder scratch memory to 32 MiB llama/30358 llama/31062
• Reordered-quantized weights: fix GET_ROWS, fix prompt-processing regression, support reordered Q8_0 MoE weights llama/30430 llama/30374 llama/29337
• Flash attention: explicit dependencies instead of host waits in the MKL path, fix crash on mismatched head dimensions llama/30684 llama/30286 llama/30812
• Fix CONV_2D results for non-contiguous input and an IQ3 crash llama/30363 llama/30604

OpenCL

• Initial Q2_K and Q3_K support across gemv/gemm/mul-mat/mul-id kernels, including MoE variants llama/30999 llama/31172
• New TOP_K and lightning indexer kernels for sparse attention llama/31026
• New ops: depthwise CONV_2D, IM2COL_3D, ARGMAX, i32 GET_ROWS and i32/f32 CPY llama/30900 llama/30895 llama/30897 llama/30756
• More DP4A GEMM coverage: cooperative kernels for Q4_0/Q4_K/Q6_K, vectorized activation loads, extra small-batch variants [llama/30741](https://github.com/ggml-org/llama.cpp/pull...

Overview

ggml v0.27.0 extends the backend scheduler API with user callbacks for split weight copying and compute-buffer reallocation, and bumps the RPC protocol to v8 to support tensor-parallel RPC servers. Backends gained substantial new fast paths: chunked gated-delta-net kernels on CUDA/HIP, KleidiAI SME2 k-quant matmuls and AVX2 Q2_0 dot products on CPU, L2-aware coopmat2 matmul scheduling and bf16 activations on Vulkan, reordered-layout MoE kernels (MXFP4, IQ3, Q5_K, grouped XMX GEMM) on SYCL, dp4a MoE GEMMs on OpenCL, tensor-API and wider-head-dim flash attention on Metal, and a DMA/HVX copy overhaul with new pooling and 3D im2col ops on Hexagon. Correctness work concentrated on non-contiguous and strided views, grid/workgroup size limits, out-of-bounds accesses, and well-defined mul_mat_id behavior for duplicate expert ids.

API changes

• Add a scheduler copy callback (ggml_backend_sched_set_copy_callback) for user-side copying of split weights (llama/29943)
• Add a scheduler realloc callback (ggml_backend_sched_set_realloc_callback) for compute buffer reallocations (llama/30329)
• Bump RPC_PROTO_MAJOR_VERSION to 8 for the new tensor-parallel RPC server mode (llama/26610)
• Document that ggml_mul_mat_id is undefined for ids that repeat an expert for the same token (llama/30501)
• Document that ggml_gated_delta_net only guarantees the latest state slot when n_tokens > K (llama/29353)
• Default symbol visibility is now hidden, removing internal symbols accidentally exported through GGML_API (llama/30014)

Core changes

• Fix out-of-bounds write in ggml_acc/ggml_set when a negative offset sign-extends past the bounds check (llama/30135)
• Fix ggml_conv_1d result layout for batched input so the output is [N, OC, OL] (llama/28738)
• Make CLAMP correct on non-contiguous views in the CPU and CUDA backends (llama/29517)
• Meta backend: treat views of host-buffer tensors as unsplit and skip them in subgraph mapping (llama/30217)
• Meta backend: recreate all subgraphs when the compute arena is reallocated (llama/30453)
• RPC: make GGML_RPC_DEBUG a verbosity level, add transport logs and clearer tensor transfer logs (llama/29544, llama/30481)

Backend changes

CPU

• Add AVX2 dot product kernel for Q2_0 (llama/30447)
• KleidiAI: add SME2 matmul for Q4_K and route Q6_K to the Q8_0 path (llama/29219)
• Restore the FP16->FP32 lookup table when F16C is available, and fix AVX2 builds not using F16C (llama/30342, llama/30343, llama/30270)
• Vectorize the fp32->fp16 conversion and optimize pad_f32 (llama/30157, llama/30526)
• Add F16 reference implementation and validation for SWIGLU_OAI (llama/29200)
• Size the set_rows conversion buffer for all threads (llama/30454)

CUDA / HIP

• Add chunked gated-delta-net kernels, plus more GDN state columns per warp (llama/29353, llama/30087)
• Support arbitrary striding in unary ops on f16/f32/bf16 and byte-strided ROLL on non-contiguous inputs (llama/29781, llama/29547)
• Improve top-k algorithm selection (llama/28713)
• Add BF16 support for XIELU and chunk BF16/FP16->FP32 conversions (llama/29955, llama/29442)
• FWHT for block widths above 512, SIMT implicit GEMM conv2d without WMMA, better pool2d coalescing (llama/29100, llama/30325, llama/29425)
• Prefer MMVQ for MUL_MAT_ID on sm_60 and pass src1 precision to host MMQ config helpers (llama/27828, llama/30168)
• Fix MMQ out-of-bounds reads and alignment assumptions when bias-add is fused (llama/29953, llama/30133)
• Handle tensors beyond grid dimension limits in the PAD and norm family kernels (llama/30147, llama/28175)
• Fix flash attention race when combining partial results and drop redundant copies after SSM_SCAN (llama/30103, llama/29807)
• Fix the CCCL version guard on major version rollover and CPU/GPU rounding agreement under MSVC (llama/29453, llama/30229)
• Add GGML_CUDA_GRAPH_DEBUG, read backend env vars once, and use per-thread streams for buffer-init memset (llama/30327, llama/28782)
• HIP: 64-row MMQ tiles for selected quant types on RDNA3.5 (llama/28195)
• MUSA: use the tile lightning indexer kernel, fix FWHT, drop an obsolete default architecture (llama/30080, llama/30167, llama/30203)

Metal

• Add a BF16 tensor-API flash attention kernel and new 128/96 head-dim FA kernels combinable across dtypes (llama/30491, llama/30209)
• Extend the few-row MMA mat-mul to the remaining src0 types and fix tensor-API mul_mm on large tensors (llama/30065, llama/28748)
• Fix MUL_MAT+ADD fusion when the residual is itself a MUL_MAT, and excess threadgroup memory in quantized FA (llama/30100, llama/29340)
• Fix mul_mat_id with duplicate expert ids (llama/30373)

Vulkan

• L2-friendly workgroup scheduling for coopmat2 matmuls, and dispatch splitting beyond workgroup count limits (llama/30215, llama/29533)
• Support bf16 activations in mul_mat and extend sparse flash attention to coopmat2 (llama/29069, llama/30003)
• Handle duplicate mul_mat_id ids in the prepass instead of looping, and use 4 rows for NVIDIA MUL_MAT_ID MMVQ (llama/29998, llama/29274)
• Overhaul host-visible memory so several devices can share it, and fix concurrent first-device initialization (llama/29741, llama/29766)
• Copy strided f16 KV once on AMD so reads spread across memory channels (llama/27703)
• Pack FMAs for f16 accumulation without the dot product op, and speed up large-window pool_2d (llama/29877, llama/30119)
• Fixes: TOP_K with +inf/NaN and k = 1 on negatives, rms_norm workgroup overflow, slow AMD iGPU reads (llama/30107, llama/30145, llama/30049)
• Intel GPU workarounds: handle free memory reporting above total memory, set shader counts for Arc Pro B65/B70 (llama/29835, llama/30431)

SYCL

• Accelerate MoE with arithmetic decoding and weight reordering for MXFP4, IQ3 and Q5_K reorder layouts (llama/29809, llama/29107, llama/29500, llama/29375)
• Add grouped MoE XMX GEMM and fuse the delta-net alpha gate (add + unary + mul) ([llama/29245](ggml-org/llama.cpp#2...

Overview

ggml v0.26.0 brings a new alloc_buffer_n/get_alloc_size_n buffer allocation API, sparse flash attention kernels on SYCL, Vulkan and Metal, and major lightning indexer improvements (halved score memory, tiling, MUSA support). The CPU backend gains BF16 ops and a tiled k-quant mul_mat, CUDA gains a model-driven W4A4 (NVFP4/MXFP4) mul_mat path plus MMVQ shared-expert fusion, and the Hexagon backend adds a sampler and more quant types. Model loading is faster with stricter GGUF size validation, Windows ARM64 MSVC builds are enabled, and WebGPU/OpenVINO/OpenCL/SYCL pick up numerous new kernels, ops and fixes.

API changes

• Add alloc_buffer_n and get_alloc_size_n to the buffer type interface with public APIs ggml_backend_buft_alloc_buffer_n/ggml_backend_buft_get_alloc_size_n for multi-buffer allocation and size planning (llama/23671)
• Add new public header include/ggml-zdnn.h for the zDNN backend (llama/29541)

Core changes

• Speed up model loading and fix a hang on crafted GGUFs with very large KV dimensions (llama/29598)
• Harden tensor/GGUF size validation: fix integer overflows and reject tensor sizes that wrap after padding (llama/29384, llama/26979)
• Collect all input tensors into graph_inputs up front and require them to be GGML_OP_NONE, fixing spurious graph re-reserves with pipeline parallelism (llama/29634, llama/29647)
• Meta backend: clear inactive AllReduce shards with FILL instead of SCALE (llama/29793)
• ggml-quants: avoid invalid rounding in the qkx3 scale search (llama/29817)
• Enable Windows ARM64 builds with MSVC cl.exe (llama/28362)

Backend changes

CPU

• Support BF16/FP16/FP32 K tails in tinyBLAS on x86 (llama/29806)
• BF16 op support: add unary, GLU, binary and scale ops (with CUDA) and accept BF16 in src1 of mul_mat (llama/29675, llama/28937)
• Tiled mul_mat for k-quants via int8 unpack tiles and 16x16 microkernels (llama/27851)
• Enable tiled flash attention for non-vector-multiple head dims on x86 (llama/29423)
• Accumulate f16 dot products in f32 on AVX512-FP16 (llama/29545)
• Add Q8_0 IME1 matrix kernel for SpacemiT X60 (llama/28479)
• Fix soft_max_back wrong output when dst aliases src1 (llama/27096)

CUDA / HIP

• Add model-driven W4A4 (NVFP4/MXFP4) mul_mat path with llama_prec_policy (llama/24364)
• NVFP4: optimize MMQ accumulation and handle the compute type on the cuBLAS path (llama/29857, llama/29173)
• Fuse shared experts into MMVQ (llama/29184)
• Use MMVF for thin f16/bf16 mul_mat at small batch size (llama/29633)
• FlashAttention: prefer whole-tile scheduling for two-stage kernels, tune fp16 tile configs for head sizes 40-112, fix 2 broken Volta cases (llama/29435, llama/26289, llama/29803)
• Lightning indexer: halve the score memory (CUDA/Metal/Vulkan) and tile over keys and tokens for 4 heads (llama/29825, llama/29901)
• Fuse RMS_NORM + SCALE into one kernel (llama/29393)
• Bitonic argsort handles rows wider than one block (llama/28957)
• Add F16 support to FWHT and a F16 kernel for CONV_2D_DW (llama/29096, llama/29064)
• SSM_SCAN: support state size 96 (llama/28717)
• Fixes: batch-independent alloc_deps check, unexpected graph reallocation in k-pool models, MMQ memory fault when n_expert >> n_ubatch, cpy transposed path corrupting non-contiguous dst (llama/29986, llama/29958, llama/29941, llama/27663)
• Make the CCCL version configurable and pin it for CI (llama/29792)
• HIP: enable the fattn-mma kernel on CDNA for dkq > 256 at large batch sizes, fixing template skip and DGX Spark misdetection (llama/28907, llama/29559, llama/29572)
• HIP: optimize packed byte subtraction (llama/29478)
• MUSA: fix PH1 (MTT S5000) operator and build issues, use the vector lightning indexer kernel (llama/29193, llama/29990)

Metal

• Flash attention: add tensor API kernel for F16 KV, split FA kernels into per-dtype libraries, optimize sparse FA (llama/29570, llama/29329, llama/29377)
• FWHT: support block widths above 512 with perf optimizations (llama/29095, llama/29602)
• mul_mat: use bf16 math for mxfp4, add few-row MMA mat-mul (llama/29770, llama/29869)
• Support left and circular padding in GGML_OP_PAD (llama/29561)
• Fix graph capture and handle empty graphs, release temporary private transfer buffers (llama/29390, llama/29777)

Vulkan

• int8 coopmat1 matmul implementation for AMD RDNA3 and RDNA4 (llama/27952)
• Sparse flash attention for quantized K/V (llama/29639)
• Perf: reuse descriptor sets when bindings are constant, fuse the SCALE→SIGMOID→SCALE→HC_POST chain (llama/29280, llama/29520)
• Flash attention: fix shmem write out of bounds and stale prealloc_y reuse across FA and soft_max (llama/29988, llama/29591)
• Fix wrong results when a mul_mat reads a slice of a larger cache (llama/28956)
• Tuning: GDN kernel on Intel, RDNA4 mat_vec, disable large matmul tile on Samsung GPUs with 32KB shared memory, argsort kernel selection on Adreno (llama/29476, llama/29934, llama/28531, llama/29469)
• Fix build with legacy GLSLC without cooperativeMatrix API (llama/29409)

WebGPU

• Add MMVQ support for Q1_0/Q5_0/Q5_1/Q3_K/Q5_K/Q6_K/MXFP4 (llama/29483)
• Add bfloat16 support for MUL_MAT/MUL_MAT_ID/GET_ROWS and f16 support for fill/set_rows (llama/29358, llama/29897)
• Fix SSM_SCAN binding aliasing and unaligned writes in buffer_set_tensor (llama/29750, llama/29471)

SYCL

• Add sparse flash attention (llama/28796)
• Q8_0 DMMV ESIMD and MMVQ wide load (llama/29186)
• FWHT kernels for block widths above 512, large register file for D=512 FA vec kernels (llama/29243, llama/29062)
• Avoid the slow oneDNN reference matmul and fattn (llama/28985)
• Reduce tensor allreduce sync with pinned host buffers (llama/29604)

OpenCL

• Add Q8_0 non-MoE dp4a and Q5_K gemv binary kernels (llama/29439, llama/29401)
• ...

Overview

A small maintenance release: the core graph size calculation (ggml_graph_nbytes) no longer triggers a UBSan "non-zero offset to null pointer" diagnostic, and the CI setup moves CUDA builds to the hf-jobs runner and adds a new AMD Vulkan build job.

Core changes

• Fix a UBSan error in ggml_graph_nbytes by computing the layout relative to a non-null base pointer
  (ggml#1644)

More info

• Releases and versioning of ggml-org projects

Changelog since v0.25.2

353b63b ggml : bump version to 0.25.3 (#1645)
d991c9b pi : update gg/SYSTEM.md
74b191a ggml : fix ubsan error in ggml_graph_nbytes (#1644)
58be011 ci : move cuda to hf-jobs and add amd-vulkan (#1641)

Overview

A small point release focused on backend improvements: a new CUDA convolution kernel with an implicit-GEMM fast path, Vulkan fixes and tuning (misalignment handling in convolution shaders and cooperative-matrix support for Adreno GPUs), an optimized DP4A binary kernel for OpenCL Q6_K GEMM, and a precision guard in the Hexagon backend.

Backend changes

CUDA

• added conv3d kernel with an implicit-GEMM fast path (F16 weights) and a direct fallback for F32 shapes (ggml-org/llama.cpp#29137)

Vulkan

• handle misalignment in the conv_2d and conv_3d matrix-multiply shaders (ggml-org/llama.cpp#29365)
• enable and tune VK_KHR_cooperative_matrix support for Adreno GPUs with hardware matrix cores (ggml-org/llama.cpp#29328)

OpenCL

• added optimized DP4A binary kernel for Q6_K non-MoE GEMM (ggml-org/llama.cpp#29057)

Hexagon

• reject MUL_MAT_ID when the src1 precision is F32 (not supported) (ggml-org/llama.cpp#29348)

More info

• Releases and versioning of ggml-org projects

Changelog since v0.25.1

a0d15e2 ggml : bump version to 0.25.2 (#1642)
c5bf7e2 sync : llama.cpp
10e36b6 vulkan: handle misalignment in conv_2d and conv_3d (llama/29365)
0d1fb47 vulkan: tune KHR cooperative matrix support for Adreno GPUs (llama/29328)
749a9eb sync : llama.cpp
aebfeef cuda : add conv3d with implicit GEMM (llama/29137)
9a7b3c1 hexagon: reject MUL_MAT_ID when src1 precision is F32 (llama/29348)
600ec61 opencl: add A8 Q6_K non-MoE dp4a binary kernel (llama/29057)
64302f4 scripts : make-release-desc - link previous release in changelog title (#1638)

Overview

A hotfix release focused on the CUDA backend: sparse flash attention, which had been disabled due to a batch-dependent gate, is re-enabled for long-context prefill (up to ~1.4x prefill speedup at 131k context), with the supporting sparse mask scan kernel significantly sped up. Additionally, Metal gains the missing f32 × bf16 matmul kernels (fixing depthwise convolutions with bf16 weights) and its flash-attention tuning tables are now keyed by GPU family instead of individual SKU, while Vulkan adds MMQ/MMV matmul kernels for the IQ4_XS quantization type.

Backend changes

CUDA

• Re-enable sparse flash attention for long-context prefill (switched off by llama.cpp#28770 due to a
  batch-dependent gate), with up to ~1.4x prefill speedup at 131k context (llama.cpp#29298)
• Speed up the sparse mask scan kernel via compile-time loop unrolling and host-side bound selection (586 → 244 us
  for the batched sparse op at 49k context) (llama.cpp#29298)

Metal

• Add the missing f32 × bf16 mul_mv kernel variants, fixing depthwise convolutions with bf16 weights (llama.cpp#28741)
• Key the flash-attention vector tuned tables by GPU family instead of individual SKU, falling back to the
  baseline for untuned families and substantially shrinking the tuning tables (llama.cpp#29075)

Vulkan

• Add MMQ/MMV matmul kernels for the IQ4_XS quantization type (llama.cpp#28415)

More info

• Releases and versioning of ggml-org projects

Changelog since v0.25.0

e565a8f ggml : bump version to 0.25.1 (#1637)
ef97dbf sync : llama.cpp
2b3b7c7 CUDA: add a reserve to avoid spurious warning on older GCC builds (llama/29317)
60e21b4 metal: add the missing f32 x bf16 mul_mv variants (llama/28741)
4506ab1 CUDA: enable sparse-fa for dsv4 prefill (again) (llama/29298)
a70078a metal : key the fa-vec tuned table by family instead of SKU (llama/29075)
8b93b8f vulkan: add IQ4_XS MMQ/MMV matmul kernels (llama/28415)

Overview

This release expands hyper-connection, flash-attention, and fused MoE/SSM support across CPU, GPU, and accelerator backends. It also improves backend robustness, quantization, data-layout handling, and RPC/meta buffer management. Numerous correctness fixes and performance tuning land across all major backends, including new kernels, fusions, and op coverage.

API changes

• Added gated ggml_dsv4_hc_pre_gated() (llama/28901).
• Allowed ggml_dsv4_hc_post() NULL comb (llama/28901).
• Bumped the RPC protocol major version to 7 (llama/24292).

Core changes

• Added gated hyper-connection preprocessing (llama/28901).
• Added hc_post identity path without comb (llama/28901).
• Fixed dimension/stride truncation in ggml_permute (llama/29227).
• Scheduler propagates graph reservation failures (llama/26070).
• Optimized IQ1_M with per-block prefix sums (llama/28706).
• GGUF data alignment is now relative to GGUF start (llama/28993).

Backend changes

CPU

• Added ARM repack kernels for Q1_0 (llama/23492).
• Fixed SpacemiT int16 transpose (llama/25161).
• Added F16 input for FWHT and hadamard MUL_MAT (llama/27779).

CUDA / HIP

• Added implicit-GEMM conv2d support (llama/29135).
• Convert contiguous tensors four elements at a time (llama/29155).
• Added i16/i32 DUP support (llama/28897).
• Added row-contiguous SUM_ROWS support (llama/26308).
• Refactored flash attention shared-memory swizzle (llama/28536).
• Tuned flash attention shapes for Ampere and newer (llama/29152).
• Enabled sparse flash attention (llama/28770).
• Tuned the SM70 MMVQ/MMQ crossover (llama/28912).
• Fixed the sm_70 tile compilation error (llama/29224).
• Fixed CUB argsort corruption from in-place keys (llama/28389).
• Fixed top-k MoE dispatch (llama/28432).
• Improved im2col access patterns (llama/28013).
• HIP: optimized IQ2/IQ3 with SWAR (llama/27962).
• HIP: enabled AllReduce for ROCm (llama/27825).
• HIP: broadened RDNA3.5 MoE tile heuristics (llama/28935).
• HIP: use F32 accumulation for MFMA flash attention (llama/28576).

Metal

• Added MoE and SSM_CONV fusion optimizations (llama/28948).
• Added hyper-connection fusion optimizations (llama/28948).
• Added arbitrary hc support in dsv4_hc_pre (llama/29169).
• Added new hyper-connection ops (llama/29000).
• Added F16 input support to FWHT (llama/29094).
• Added flash attention kernels for new head sizes (llama/28599).
• Fixed flash attention mask bounds (llama/29220).
• Fixed flash attention support checks (llama/29122).
• Fixed mul_mm_id NaN for large F16 activations (llama/26223).
• Gated mul_mm_id src1 rescale behind ggml_prec (llama/29029).
• Fixed macOS SDK deprecation warnings (llama/29136).
• Simplified fusion-pattern declarations (llama/29206).

Vulkan

• Added IQ3_S MMQ matmul kernels (llama/28822).
• Added sparse flash attention support (llama/28105).
• Added Intel Xe flash attention optimization kernels (llama/24406).
• Added new hyper-connection ops (llama/28988).
• Refactored Vulkan buffers/debug into shared modules (llama/28732).
• Raised mul_mat_id row-id limit to 512 experts (llama/28501).
• Skipped unneeded MoE work in the coopmat1 path (llama/25483).
• Fixed buffer reference alignment in im2col shaders (llama/28996).
• Worked around the NV argsort_large compiler bug (llama/28975).
• Made MUL_MAT_ID BN/2 tail unconditional (llama/28923).

OpenCL

• Added an A8 Q4_0 non-MoE dp4a binary kernel (llama/29055).
• Added an A8 q4_K/q8_1 binary GEMM kernel (llama/29056).
• Added an A8 q6_K binary GEMM kernel (llama/28678).
• Added a binary flash attention kernel (llama/29046).
• Added generic ssm_scan support (llama/28881).
• Choose the MoE expert matmul by batch size (llama/27637).

SYCL

• Added GET_ROWS_BACK support for FP32/FP16 (llama/25266).
• Added gated dsv4_hc_pre support (llama/29132).
• Added optional hc_post comb matrix support (llama/29132).
• Added MMVQ GLU fusion (llama/28931).
• Added rms_norm+scale and ssm_conv+silu fusions (llama/28931).
• Fused the SiLU epilogue into the ssm_conv kernel (llama/28929).
• Improved MKL-FA softmax load coalescing (llama/28918).
• Fixed pinned memory to use the right device context (llama/28895).
• Fixed a large-allocation memory error (llama/28953).
• Fixed the split buffer type function signature (llama/28981).
• Added hadamard fp16 unit-test support (llama/29218).

Hexagon

• Added an HMX-optimized gated delta net (llama/29199).
• Added Q4_K and Q6_K K-quants support (llama/28994).
• Added ROLL op support (llama/29105).
• Added TOP_K op support (llama/29113).
• Added GEGLU_QUICK op support (llama/29114).
• Added I32 GET_ROWS support (llama/29116).
• Overhauled buffer/DMA handling for 64-bit mappings (llama/29197).
• Added direct-mapped DMA cache for FA masks (llama/29282).
• Updated im2col support (llama/29103).
• Added flash attention head-dim padding (llama/26539).
• Restored contiguous fast-paths and hvx_copy_uu (llama/28886).
• Use DMA for contiguous hex-cpy (llama/28906).
• Accept the zeroed rope probe in supports_op (llama/28995).

WebGPU

• Added fused gated-delta-net + copy (llama/28976).
• Fixed GET_ROWS support checks (llama/28978).
• Added self-hosted WebGPU CI coverage (llama/28712).

OpenVINO

• Updated OpenVINO to 2026.4 (llama/29009).
• Fixed clangd and MSVC warnings (llama/29009).
• Optimized stateful decode and GPU MoE inference ([llama/286...

Overview

This release focuses on expanding backend coverage and robustness, with a new precision-control API, major Vulkan/SYCL/Hexagon/OpenCL work, and numerous correctness and performance fixes across CPU, CUDA, Metal, and other backends.

API changes

• Expanded ggml_prec with GGML_PREC_BF16, F16, Q8, and Q4, and deprecated GGML_PREC_DEFAULT (llama/26675).
• Added ggml_prec_set_acc() and ggml_prec_set_src() to control accumulator and per-source precision for MUL_MAT/MUL_MAT_ID and flash attention (llama/26675).
• Deprecated ggml_mul_mat_set_prec() and ggml_flash_attn_ext_set_prec() in favor of the new precision API (llama/26675).

Core changes

• Added precision op-params layout for MUL_MAT and MUL_MAT_ID in ggml-impl.h (llama/26675).
• Backend scheduler no longer forces an extra split for backend inputs and skips zero-sized MoE ids tensors (llama/28387, llama/28739).
• Added PCH/unity-build support to speed up builds (llama/28091).
• Fixed MSVC+Clang ggml_vld1q_u32 and added a missing header for gguf.cpp (llama/28284, llama/28566).

Backend changes

CPU

• Disabled PCH and fixed CACHE_LINE_SIZE ambiguity causing heap corruption (llama/28882).
• Added s390x repack support for q4_0 and Q1_0 vector intrinsics; guarded VXE-only helpers and added non-VXE tests (llama/28667, llama/28606, llama/28775, llama/28776).
• Added PCH support to ggml-cpu (llama/28091).
• Fixed MSVC+Clang ggml_vld1q_u32 (llama/28284).

CUDA / HIP / MUSA

• Added AMD GCN-specific MMQ config table and gfx90c HIP support (llama/27841, llama/26454).
• Fall back to F32 on devices without BF16 hardware acceleration (llama/28846).
• Reworked flash attention quant compile flags (GGML_FA_QUANTS) and tuned flash attention for gfx1201 (llama/28079, llama/28102).
• Fixed divergent barrier in f16 flash attention and races in MoE MMID/MMF kernels (llama/27870, llama/28475).
• Tuned MoE MMQ N-tile sizing for RDNA3 and added branchless Q4_K/Q5_K unpack with L2 prefetch (llama/24546, llama/28552, llama/26705).
• Reverted HIP prop.integrated restore (llama/28604).

Metal

• Reworked fusion into a single-source table with debug support (llama/28164).
• Added flash attention vec tuning for M2 Max and M3 (llama/28458, llama/28396).
• Fixed idle threads in IQ mul_mv kernels and skipped empty MoE token tiles (llama/28692, llama/28086, llama/28301).
• Fixed memory leak in an early-return path (llama/28399).

Vulkan

• Added TQ1_0 and iq4_xs support, type-aligned GET_ROWS, and f16 B-type matmul pipelines (llama/27765, llama/28426, llama/28253, llama/27471).
• Added hyper-connection fused ops and RMSNorm/UNARY fusion opportunities (llama/26578, llama/28024, llama/27220).
• Added command-buffer debug labels and improved allocation dependencies for prefill topk_moe fusion (llama/28101, llama/28422).
• Fixed argsort data race/OOB, NV queue-submit driver workaround, FILL workgroup limits, and PowerVR dmmv fallback (llama/28705, llama/28830, llama/28592, llama/28341).
• Use spec constants for MM A-type and added small-M optimizations (llama/25773, llama/28457).
• Fixed 32-bit Vulkan-Hpp handle usage and added CPU write path for async copy when idle (llama/22892, llama/28618).

OpenCL

• Added Adreno xmem SDPA path and extended elementwise/data-movement op coverage (llama/26331, llama/27633).
• Added A8 q4_0/q4_K binary GEMM kernels and fixed noshuffle row alignment for q4_K/q5_K/q8_0 (llama/28268, llama/28677, llama/28575).
• Fixed backend abort bugs, non-contiguous conv2d inputs, and q4_K/q5_K weight pack selection (llama/27630, llama/28503, llama/28402).

SYCL

• Added radix-select top_k, batched L2_NORM, and IQ type handling for MoE (llama/28670, llama/28222, llama/28476).
• Restored Kronecker product FWHT support and added allocation site tracing (GGML_SYCL_MEMTRACE) (llama/28254, llama/27631).
• Fixed oneDNN scratchpad pool ordering and memory error handling (llama/28704, llama/28227).
• Fail gracefully for unsupported TQ1_0 quants (llama/28681).

Hexagon

• Added multi-device model split (row-split) support and rope updates (llama/28589, llama/28628).
• Added RELU and LEAKY_RELU ops (llama/28585).

WebGPU

• Updated Dawn and aligned tensor bindings to type block size (llama/28683, llama/28382).
• Added type-aligned GET_ROWS support (llama/28253).

RPC / OpenVINO

• Fixed RPC linking with BUILD_SHARED_LIBS=OFF (llama/28492).
• Updated OpenVINO for type-aligned GET_ROWS (llama/28253).

More info

• Releases and versioning of ggml-org projects

Changelog since v0.23.0

456172e ggml : bump version to 0.24.0 (#1627)
d71da1d sync : llama.cpp
869a1fe tests(s390x): add non-vxe build to tests (llama/28776)
cc40f25 sycl: rfc: Use radix select for top_k (llama/28670)
5eccfef ggml-cpu : disable PCH and fix CACHE_LINE_SIZE ambiguity to fix heap corruption (llama/28882)
9613d22 sycl : fix oneDNN scratchpad breaking the pool free order (llama/28704)
c250e9c ggml-cuda: fallback to F32 on device without BF16 hardware acceleration (llama/28846)
0a3eaa5 sync : llama.cpp
bad30a1 ggml-cpu(s390x): guard VXE-only repack helpers (llama/28775)
80fc34b sycl : Fix get mem error (llama/28227)
ce73076 tests : reduce FA test sizes (llama/28842)
2e3a01e vulkan: workaround NV queuesubmit driver bug (llama/28830)
4c0b472 opencl: apply the noshuffle row-alignment rule to q4_K, q5_K and q8_0, not just q6_K (llama/28575)
d214dee sync : llama.cpp
a389817 ggml-cuda: hip add specific config table for AMD GCN (llama/27841)
93424a0 syscl : Handle (fail gracefully) unsupported tq1_0 quants (llama/28681)
807b5c3 rpc : fix linking when compiling with BUILD_SHARED_LIBS=OFF (llama/28492)
e473e15 opencl: fix several bugs where the backend aborts (llama/27630)
be4ef89 opencl: add bin kernel kernel_gemm_noshuffle_q4_k_f32_32b_trans_ila_a8_bin (llama/28677)
0ccbed8 webgpu: align tenso...

Overview

Version 0.23.0 adds new operators and backend scheduling APIs for sparse attention, asynchronous execution, and allocation-dependency tracking. It includes broad correctness and performance work across CPU, GPU, and accelerator backends, plus an expanded RPC stack with Apple RDMA support.

API changes

• Added GGML_GLU_OP_SWIGLU_CLAMP and ggml_swiglu_clamp() (llama/27930).
• Added ggml_flash_attn_ext_set_n_kv_max() for sparse K/V masks (llama/27970).
• RPC protocol bumped to 6.0.0 for event/async APIs (llama/18626).

Core changes

• Added SWIGLU_CLAMP and sparse-attention n_kv_max support (llama/27930, llama/27970).
• Added allocation-size expansion tracking and allocation-dependency support in backend graph optimization (llama/27960, llama/27301).
• Fixed multi-batch conv_transpose_2d and allowed SET_ROWS in backward passes (llama/26132).
• Backend search-path enumeration tolerates filesystem errors; meta buffers propagate usage (llama/28271, llama/27586).
• Replaced compile-definition versioning with generated ggml-version.h (llama/28364).
• RPC gains event/async APIs, Apple RDMA transport, and avoids serializing remote buffers (llama/18626, llama/26421, llama/26500).

Backend changes

CPU

• AVX2 IQ prompt speedups for large batches (llama/27402).
• Fixed s390x Q5_1 uninitialized accumulator and added conditional SpacemiT IME kernels (llama/28332, llama/27961).
• Reworked KleidiAI build/integration and avoided buffer-type init on dispatch (llama/26077, llama/27891).

CUDA / ROCm

• Added sparse FA, XOR-swizzled FA K/V smem, radix top-k, and extended MoE fusions (llama/27970, llama/25635, llama/27466, llama/25952).
• Added concurrent streams per split, MMQ/MMVQ tunings, and MoE MMQ fixes (llama/28198, llama/28285, llama/26264).
• ROCm/hip: tuned RDNA3 MMQ configs and optimized Q2_0 dot products (llama/26284, llama/26753).

Metal

• Added sparse FA, radix top-k, quantized concat, and chunked SSD MMA (llama/28098, llama/28073, llama/28116, llama/26647).
• Added fa-vec tunings across many Apple GPUs and Metal 4.0 tensor API support (llama/28373, llama/27461).
• Added metallib xcframework support and fixed memory, leak, and dispatch issues (llama/28163, llama/27758, llama/28306).

Vulkan

• Added cross-entropy loss, radix top-k, Lightning Indexer op, and IQ3_S batch mat-vec improvements (llama/27216, llama/28032, llama/27453, llama/27449).
• Fixed FA dequant, mul_mat_id padding, view-alias deps, and warp-size handling (llama/28190, llama/27925, llama/27812, llama/27726).
• Added shader hoisting for row IDs/expert counts and bfloat16 extension gating (llama/26686, llama/28155).

SYCL

• Added residual fusions, peer-to-peer copy, and host-pinned memory limits (llama/27610, llama/27550, llama/27559).
• Added GPU free-memory query, split TOP_K long rows, TILE quantized KV decode, and oneDNN SDPA f16 KV binding (llama/27968, llama/27847, llama/26689, llama/27468).
• Reduced redundant Q4_K MMVQ work and improved --fit handling (llama/27062, llama/27629).

OpenCL

• Added quantized lm_head/decode GEMV and medium-batch GEMM optimizations (llama/26477).
• Added Adreno fixes and matmul path improvements, Intel Xe-LP quant tuning, and MoE binary kernels (llama/27640, llama/26438, llama/27768).

Hexagon

• Added F16 unary ops, ABS/LOG, MUL_MAT/MUL_MAT_ID fusion, and CPY fence fixes (llama/28228, llama/27786, llama/28202, llama/28033).
• Added multi-NPU device support, fully async backend, and on-demand device discovery (llama/26501, llama/27785).

OpenVINO

• Updated OpenVINO, added new op support, whisper.cpp support, and a fuse-to-conv pass (llama/27843).

RPC

• Implemented event and async backend APIs and Apple RDMA transport (llama/18626, llama/26421).
• Avoided serializing remote buffers and fixed pre-RDMA macOS teardown/version issues (llama/26500, llama/27815, llama/27908).

WebGPU

• Fixed misaligned offset crash in tensor get and marked MUL_MAT for allocation expansion (llama/28045, llama/28071).

More info

• Releases and versioning of ggml-org projects

Changelog since v0.22.0

e91ded1 ggml : bump version to 0.23.0 (#1618)
d8e6687 sync : llama.cpp
7523a09 metal : add remaining fa-vec tunings for M3 Max (llama/28373)
72d22bd ggml : replace compile definitions with version.h.in (llama/28364)
7aced59 ggml : rename and make private ggml_op_alloc_size_may_expand() (#0)
24b5dee sync : llama.cpp
b849c94 ggml : don't crash when backend search path can't be read (llama/28271)
da8baab ggml : remove GGML_CUDA_PEER_MAX_BATCH_SIZE (llama/28177)
3573afc ggml-cpu(s390x) : fix q5_1 uninitialized v_acc (llama/28332)
b70ff92 sycl: fuse rms_norm+mul+add and add+add residual chains (llama/27610)
db62b97 SYCL: Refactor GGML_SYCL_ENABLE_MKL_FA to global var (llama/26863)
6921674 opencl: quant lm_head / decode GEMV and medium-batch GEMM optimizations (speculative decoding/MTP) (llama/26477)
90c4733 tune MMVQ to MMQ crossover for SM87 (llama/28285)
1145b17 metal : add sparse FA (llama/28098)
8ed43cf metal : fix glu dispatch with ne00 = 1 (llama/28306)
eb62444 CUDA: Allow concurrent streams per split for multi-GPU (llama/28198)
355bd37 vulkan: fix FA dequant path engagement (llama/28190)
04c25ad sycl : enhance the api to support peer-to-peer copy (llama/27550)
db157a7 sycl: reduce redundant work in Q4_K multi-column MMVQ (llama/27062)
ee9842c finetune: fix no KV cache (llama/27199)
8d6211c ggml-hexagon: add F16 support for unary ops (llama/28228)
1e7de5d metal : add fa-vec tunings for M3 (llama/28236)
f78f88b metal : fix memory query under low-memory conditions (llama/27701)
b40e3c4 ggml-cuda : remove unused vars (llama/28235)
589cf31 CUDA + ggml: add sparse-fa for DSV4/GLM (llama/27970)
b9ba0ae ggml: avoid KleidiAI buffer type init on dispatch (llama/27891)
d9d16fd hexagon: MUL_MAT and MUL_MAT_ID fusion and fixes (llama/28202)
a4325bb vulkan: handle larger batch sizes (>4) effi...