ggml Files

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

456172ec ggml : bump version to 0.24.0 (#1627) d71da1df sync : llama.cpp 869a1fea tests(s390x): add non-vxe build to tests (llama/28776) cc40f25e sycl: rfc: Use radix select for top_k (llama/28670) 5eccfef7 ggml-cpu : disable PCH and fix CACHE_LINE_SIZE ambiguity to fix heap corruption (llama/28882) 9613d224 sycl : fix oneDNN scratchpad breaking the pool free order (llama/28704) c250e9cf ggml-cuda: fallback to F32 on device without BF16 hardware acceleration (llama/28846) 0a3eaa56 sync : llama.cpp bad30a14 ggml-cpu(s390x): guard VXE-only repack helpers (llama/28775) 80fc34b9 sycl : Fix get mem error (llama/28227) ce730762 tests : reduce FA test sizes (llama/28842) 2e3a01e1 vulkan: workaround NV queuesubmit driver bug (llama/28830) 4c0b4724 opencl: apply the noshuffle row-alignment rule to q4_K, q5_K and q8_0, not just q6_K (llama/28575) d214dee9 sync : llama.cpp a3898179 ggml-cuda: hip add specific config table for AMD GCN (llama/27841) 93424a02 syscl : Handle (fail gracefully) unsupported tq1_0 quants (llama/28681) 807b5c36 rpc : fix linking when compiling with BUILD_SHARED_LIBS=OFF (llama/28492) e473e155 opencl: fix several bugs where the backend aborts (llama/27630) be4ef895 opencl: add bin kernel kernel_gemm_noshuffle_q4_k_f32_32b_trans_ila_a8_bin (llama/28677) 0ccbed89 webgpu: align tensor bindings to the type block size (llama/28382) 000ca388 hexagon: support for multi-device model split (aka row-split) (llama/28589) 01f6b9f6 ggml-webgpu: Update to a recent version of Dawn (llama/28683) 19037b49 tests: extend test-quantize-fns to test nrc=2 (i8mm) kernels (llama/16234) 1c54744b ggml: skip 0-sized ids tensor when offloading selected experts (llama/28739) bdc63ed0 metal : skip the empty half of the mul_mm_id token tile (llama/28301) 47ddafe5 cmake : add PCH and unity build to improve build times (llama/28091) 456a1b8c metal : single-source fusion table + fusion debug rework (llama/28164) e44fc424 metal : fix idle threads in the remaining iq mul_mv kernels for ne00 < 1024 (llama/28692) c69da0a9 CUDA/HIP: Flash Attention tuning (gfx1201) (llama/28102) 2e2531e2 vulkan: fix data race and OOB access in argsort(large) (llama/28705) 7ecf3fd7 opencl: add A8 Q4_0 mm binary kernel support (llama/28268) d124cfb5 vulkan: use CPU writes in ggml_backend_vk_cpy_tensor_async if the context is idle (llama/28618) 33abc912 vulkan: use add_alloc_dep to enable topk_moe fusion for prefill (llama/28422) 83ecb318 vulkan: small M matrix optimizations for qwen (llama/28457) 68e8fdae tests : increase tolerance for Add fusion tests (llama/28691) d93b5bce tests : drop SYCL special-casing in test-backend-ops.cpp (llama/28688) c35fdb2e vulkan: fall back to shared-memory reduction for dmmv on PowerVR (llama/28341) 6551d151 vulkan : add command-buffer debug labels for GPU profilers (llama/28101) 678d199b ggml-cpu(s390x): add repack support for q4_0 (llama/28667) 0e260d14 ggml-cpu(s390x): add Q1_0 vector intrinsic support (llama/28606) a37fc514 vulkan: use spec constant for matrix matrix multiplication A-type (llama/25773) 3b970179 hexagon: rope updates (llama/28628) 9fef7bd8 vulkan: Convert FILL to distribute workgroups in 2D to avoid exceeding maxComputeWorkGroupCount (llama/28592) 7840aaba sync : llama.cpp 78a19910 CUDA: size routed MoE MMQ N-tiles from typical expert width on RDNA3 (llama/28552) f6996d7a CUDA: replace GGML_FA_ALL_QUANTS with GGML_FA_QUANTS, more control over what is compiled (llama/28079) 44323b47 vulkan: add dedicated iq4_xs mat-vec shader (llama/28426) 4bbe406c vulkan: add f16 B-type matmul pipelines and warp tile size tuning for Intel coopmat1 (llama/27471) fe2f7fd8 tests : use less threads for data initialization (llama/28325) c45266b2 Add IQ type handling for MoE (llama/28476) 1d992e29 ggml : fix msvc+clang ggml_vld1q_u32 (llama/28284) 2146233a Revert "ggml-cuda : restore prop.integrated on HIP builds (llama/24233)" (llama/28604) 418e0208 metal : fix idle threads in mul_mv_iq3_xxs for ne00 < 1024 (llama/28086) ef036b11 llama : add missing headers (llama/28566) 1a56192c vulkan : fuse UNARY(GELU|SIGMOID|SILU|SOFTPLUS) + MUL (llama/27220) ca0c16a8 Fix Vulkan-Hpp handle usage on 32-bit targets. (llama/22892) d762adad opencl: properly handle non-contiguous inputs to conv2d (llama/28503) 6639870e ggml : update ggml_prec specification (llama/26675) 2979e4b9 hexagon: add RELU and LEAKY_RELU ops (llama/28585) 07e28d5d tests : initialize the L2_NORM batch array (llama/28553) e39043bc Revert "CUDA: size routed MoE MMQ N-tiles from typical expert width on RDNA3 (#24546)" (llama/28551) 05adedd4 webgpu: format the GET_ROWS case block (llama/28542) 699240a6 sycl: add a batched L2_NORM kernel (llama/28222) af1e2ff8 vulkan: add DeepSeek-V4 hyper-connection fused ops (DSV4_HC_COMB/PRE/POST) (llama/26578) 886b12b6 CUDA: size routed MoE MMQ N-tiles from typical expert width on RDNA3 (llama/24546) a9d11292 ggml: add gfx90c HIP support (llama/26454) 4aced562 CUDA: branchless Q4_K/Q5_K unpack to speed up mmvq, L2 prefetch on DGX Spark (llama/26705) 04d63b27 vulkan: support type-aligned GET_ROWS (llama/28253) 39e54755 ggml-cuda: fix divergent barrier in f16 flash attention (llama/27870) eada51f2 ggml: allow backend inputs to not create another split (llama/28387) bdd7503f vulkan: rms_norm fusion opportunities (llama/28024) 7eb8fcf2 vulkan: add TQ1_0 support (mm, mat-vec, mat-vec-id, dequant, get_rows) (llama/27765) 33ad7da0 opencl: properly choose weights pack for q4_K, q5_K mul_mat (llama/28402) 7ae53f95 cuda: fixes races in mmid and mmf (llama/28475) f2881302 metal : add remaining fa-vec tunings for M2 Max (llama/28458) 85c09196 metal : fix memory leak in early return (llama/28399) 0d6e721b sycl : fix test-backend-ops CI break && restore Kronecker product FWHT support (#28016) (llama/28254) 1ac03c74 sycl: attribute device allocations by site (GGML_SYCL_MEMTRACE) (llama/27631) d7ff702c metal : add remaining fa-vec tunings for M3 (llama/28396) 23881312 opencl: extend the elementwise and data‐movement op coverage (#27633) ddddc180 opencl: add Adreno xmem SDPA path (llama/26331)