ggml Files

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 llama/30753 llama/30850
• Precompiled MxFP4 MoE DP4A prefill GEMM kernel for Adreno X1 llama/30691
• Write KV-cache views in set_tensor so state restores are not silently dropped llama/30471

WebGPU

• Add Q2_0 support and handle the -1 expert index in the MoE matmul llama/30914 llama/26631

Hexagon (HTP)

• Rework tensor repacking for tiled and flat quantized and floating-point weights, now vectorized and threaded llama/30635 llama/30901
• New Q1_0 and Q2_0 weights, with MUL_MAT using Hadamard rotation and GET_ROWS for flat and tiled layouts llama/30122 llama/30627
• Flash attention: KV-head team split for multicore HMX, multi-stream execution, long-context and sliding-window fixes llama/31055 llama/30218 llama/30731 llama/30581
• Op fusion reworked, adding new fusions such as chained elementwise ops llama/30843
• Generic dimension collapse/reshape for unary, binary and matmul ops, and row skipping in the MoE matmul llama/31071 llama/31169
• New ops SGN and CEIL, plus HVX-vectorized SOFTPLUS llama/30346 llama/31006 llama/30435
• Configurable VTCM size, fix VTCM size query for graphs with layout out-params, toolchain v0.8 with Hexagon v85 llama/31221 llama/30655 llama/30854

ZDNN

• Introduce flash attention (FLASH_ATTN_EXT) support llama/30512

CANN

• Fix #ifdef blocks that left 310P-only checks falling through into unsupported paths, add missing const qualifiers llama/25273 llama/30933

More info

• Releases and versioning of ggml-org projects

Changelog since v0.27.0

facdc44a ggml : bump version to v.28.0 (#1676) 23261912 sync : llama.cpp 5e5175f3 cuda : add top-k MoE fusion for 384 experts (llama/31288) 3f866f9b vulkan: clamp A prefetch to end_k in coopmat1 int8 matmul (llama/30283) 06db896e HIP/CUDA: remove pragma unroll from loops in conv2d where itteration count cant be infered at compile time (llama/31255) 85dcbfe4 rpc : catch std::length_error in recv_msg (llama/31204) d1989dd2 vulkan: fix CI llvmpipe by allowing dim 3 broadcast in MUL_MAT (llama/31239) 0981f4fd metal : add GGML_OP_TOP_K_MASK (llama/31230) 73cdd6f7 ggml : relax scale to non-contiguous views (llama/31146) 0e91cde0 opencl: refine mul mat id -1 index (llama/31170) 026a2eb3 opencl: initial support for q3_k (llama/31172) a9361435 hexagon: add GGML_HEXAGON_VTCM_SIZE (llama/31221) 687b158a opencl: write KV-cache views in set_tensor so a state restore is not dropped (llama/30471) 8abd879a hexagon: MUL_MAT_ID row skipping optimizations and improvements (llama/31169) 8284b4b2 ggml-cpu : rewrite concat as one row kernel for all types (llama/31179) 7abd2f26 sync : llama.cpp 865631d0 vulkan: rms norm optimization using subgroup reductions (llama/29882) d1ba9f36 sync : llama.cpp 9e4cb3cc sycl: template the plain MMVQ launch triples (llama/31013) b5186728 sycl: template the per-quant concat impls (llama/31008) 93133a7d sycl: add set_tensor_2d for faster tensor-split loading (llama/30358) 6afa37e6 sycl : bound reorder temp memory to 32 MiB (llama/31062) 4542ff2e ggml-cpu : fix unused variables in arm arch (llama/31180) 10bcc3f3 ci : fix race in webgpu memset (llama/31162) 2be37268 metal : stop overflow from happening in large matrix output offsets (llama/30072) 1b31c050 deepseekv4.1: follow-up review fixes (llama/31137) e9fc24b9 cuda : use std::string instead of VLA for archName (llama/30463) 1899b2ef CUDA: fuse FFN gate/up matmuls and GLU in MMQ with interleaved gate/up rows (llama/29948) 94506491 ggml-cuda: fuse qwen MRoPE + set-rows (llama/31025) e1599525 vulkan: Add a larger tile size for conv3d (coopmat2 only) (llama/29920) 8d226fe9 vulkan : one warp tile per subgroup in the small matmul tiles (llama/30106) 592a906a hexagon: generic dims collapse/reshape for unary, binary and matmul ops (llama/31071) 6706d832 opencl: add topk and lightning indexer (llama/31026) 37d700db opencl: add initial q2_k support (llama/30999) 4212b0d6 model: DSV4.1 flash (llama/31086) 06984d35 hexagon: KV-head team split for multicore HMX flash attention (llama/31055) d22f1156 metal : remove mul_mm_id amax rescaling (llama/30924) 45b4b384 rpc : add bounds check in rdma_recv (llama/30399) 843436ec llama : deny -sm tensor for arches with moe amax rescaling (llama/31081) 568ba065 ggml : allow MUL_MAT_ID -1 index to skip the computation (llama/26631) edc03add vulkan: ggml src1 precision support (llama/30656) 88071185 HIP: add byte permute fallback when __builtin_amdgcn_perm is unavailable (llama/30915) 1fd478aa hexagon: add SGN op (llama/30346) b3844edf hip: fix bf16 WMMA operand vector type for gfx11/gfx12 (llama/30885) 59708242 ggm-cuda: Enforce DP4A MMQ tile height multiple of warp size (llama/30837) 752716f5 CUDA: fix build after chunked CUDA graphs (#30804) (llama/31053) 8a6e170a qwen4exp : build the QSA selection without the kq_mask (llama/30965) 310a27aa CUDA: chunk graphs to amortize graph launch (llama/30804) b217b632 hexagon: add CEIL op (llama/31006) 461c2128 ggml : fix time overflow in ggml_time_us (llama/30969) b44c81b2 CANN: missing const qualifiers added for consistency (llama/30933) 6c7b935e sycl : pass device into mul_mat_uses_onednn (llama/31014) 74ecf131 sync : llama.cpp 9ce09590 hexagon: Op Fusion refactoring, optimizations and new fusions like BINARY_NX (back-to-back ADDs, etc) (llama/30843) 6ea1993f sync : llama.cpp c64f2ec5 sycl: Replace host waits in fattn-mkl with explicit dependencies (llama/30684) fdd4e61e sycl: Add graph record and replay (llama/28725) 4b72143a opencl: add depthwise conv2d (llama/30900) ea93f506 Multi-column on Intel XMX matrix engines refactored (llama/29864) e05d3b3f vulkan: replace std::aligned_alloc with posix_memalign for termux compatibility (llama/30828) 3c02b215 sycl: fix CONV_2D wrong results for non-contiguous input (llama/30363) f7086816 kleidiai : add Q8_0 SVE dotprod GEMV and SVE i8mm GEMM kernels (llama/30808) 393cc292 sycl: merge the fp16 and fp32 dequant switches (llama/30740) 0e0b6be1 webgpu : add Q2_0 support (llama/30914) 5eb05063 opencl: add F32-to-I32 ARGMAX (llama/30897) b81bf009 llama: create new fused op for QSA mask (llama/30640) ed8639d2 sycl: q8 fast dequant (llama/30775) db7c146a ggml-opencl : add IM2COL_3D support (llama/30895) c47a5a8a sync : llama.cpp c15985f4 metal : fix memory leaks and cleanup logic in GPU capture path (llama/30890) 1d13c8e7 sync : llama.cpp 614decbf hexagon: neon-vectorized and threaded repack (tiling) (llama/30901) bbbf288f opencl: add dp4a GEMM variants to vector load activation (llama/30753) 496eff9d sycl : fix fa crash for Query head dimension mismatch (neq0 != HSK / D=2048/4096) (llama/30286) ec7bc58b sycl: collapse duplicate need_check launch bodies in mmq (llama/30739) 43ae36d0 ggml-cpu : compute f32 gelu_quick and geglu_quick without the fp16 table (llama/30737) 83df5951 llama : scale moe down projection per expert token (llama/30861) c60c0ed4 Sycl iq3 crash fix (llama/30604) d242ac9e opencl: fix rope program build error with newer Clang (llama/30893) 30dffe9d hexagon: fix VTCM size query on images that need layout out-params (llama/30655) 6996d1ec sync : llama.cpp 87d33b41 ggml-cuda : support ssm_conv kernel size 7 (llama/30873) 08234b3b sync : llama.cpp f6c63dda metal : optimize argmax and pool_1d (llama/30862) aaed8948 opencl: add several gemm variants for small batches and some fixes (llama/30850) c2fe0d07 sycl : replace memset due to not always ordered (llama/30812) d4f8391b ggml-cpu : select the NEON fp32->fp16 conversion (llama/30505) fb198b56 HIP/CUDA: Fix reduction in mmf getting tile widths not devisiable by 64 (llama/30836) d5ba1a7f ggml-zdnn: introduce flash attention op (llama/30512) d7d5807b CUDA: Clear CUDA graphs after legacy memory pool flush (llama/26771) 7d77885c metal : fix missing MoE kernel variant for 9 experts per token (#30610) (llama/30699) 2b9ab6bd hexagon: update to snapdragon-toolchain v0.8 and enable hexagon v85 (llama/30854) 57ea06b8 sycl: fix compiler warnings - missing prototypes, dead code, extra semicolons (llama/30271) 50e4a18e sycl : fix get_rows on reordered quantized weights (llama/30430) 6dac3408 opencl: add cooperative-k dp4a gemm for q4_0, q4_k and q6_k (llama/30741) 90e6a7a3 opencl: add get_rows for i32, cpy for i32 and f32 (llama/30756) 96135fa8 metal : fix moe_reduce fusion race when dst aliases its inputs (llama/30795) 228ecef1 cann : refactor #ifdef blocks to avoid unreachable code after return (llama/25273) 8601f639 CUDA: fix warning for nonexistent extra semicolon (llama/30774) aeadd496 ggml-cpu: vectorize ggml_v_silu and ggml_v_expf (llama/30791) aa623eaf cuda: cap MMQ Stream-K blocks at available work (llama/30553) 9be188a3 CUDA: fix shared-memory WAR hazard in flash_attn_ext_vec (llama/30573) 7cdff029 hexagon: FLASH_ATTN_EXT improvements for longer context and SWA models (llama/30731) 50df1d46 hexagon: allow multistream FA (llama/30218) db9a74c5 cuda : fuse qwen4exp's SCALE -> SIGMOID -> SCALE -> hc_post chain (llama/30602) 58103fd7 ggml: support non-contiguous SUM (CPU, CUDA, Vulkan, WebGPU) (llama/28851) e4a5c1c2 ggml-cuda: fattn do not check oob if k_VKQ_max = (k_VKQ_max - nbatch_fa) (llama/30572) 83bb0132 Restore ESIMD kernels on FFN path (llama/30239) 9bf45943 hexagon: add support for hadamard MUL_MAT and GET_ROW for flat and tiled Q1_0 and Q2_0 (llama/30627) a76bee60 hexagon: add support for Q1_0 and Q2_0 (llama/30122) d8e5e346 ggml-cuda: HIP: use __builtin_amdgcn_perm for q1_0 unpacking (llama/29927) 647863c0 CUDA: faster MMVQ for low-bit types (llama/30517) 1a7c0715 cuda : fix -Wtautological-compare in fattn vec dispatch (llama/30519) 4e9d1da0 cuda : add grouped experts top-k fusion (llama/29181) 2c22d227 sycl: fix prompt-processing regression after Q4_0/Q8_0 weight reorder (llama/30374) d347a485 hexagon: vectorize softplus with HVX (llama/30435) 3a2611b7 opencl: Add MxFP4 MoE bin kernel for Adreno X1 devices (llama/30691) 91a5a9d1 hexagon: correct HMX FA sinks (llama/30581) f287e731 fix memory errors in mul_mat, split buffer, host pool (llama/29889) 02c69c40 CUDA: HIP: rms_norm: add variants that load float4 at a time and store x in registers (llama/29720) 8fe8ade6 cpu : reduce AVX2 tiled accumulators together (llama/30296) 8c5987b7 ggml-cpu: refactor s309x repack helper functions (llama/30588) f6f84bb9 sycl: wide stores for the q4_K/q5_K weight dequant (llama/29696) 133b7e91 vulkan: fix gpu-visible host memory issues (llama/30468) 0c8d3259 sync : llama.cpp 7ce52a37 hexagon: tensor repacking rework and support for tiled and flat quantized and fp weights (llama/30635) 7415af36 sync : llama.cpp 17c3063a Limit max CUDA pinned host buffer size to 4 GB (llama/30449) ffbbe31e sycl: support reordered Q8_0 MoE weights (llama/29337) da8e2a54 ci : upload API/ABI reports as artifacts (#1666) 34a80fbe ci : fix --strict mode for API check (#1665) 43925a09 pi : turn .pi/gg into a system prompt addendum