]> git.djapps.eu Git - pkg/ggml/sources/whisper.cpp/log
pkg/ggml/sources/whisper.cpp
5 weeks agoFix data race in CUDA's "cpy" kernel (influences GGML's DUP, CONT operations). (llama...
Rail Chabdarov [Sat, 14 Mar 2026 05:19:44 +0000 (06:19 +0100)]
Fix data race in CUDA's "cpy" kernel (influences GGML's DUP, CONT operations). (llama/20507)

* Fix datarace in CUDA's "cpy" kernel.

* Remove extra barrier by using more of shared memory.

5 weeks agoopencl: fix l2_norm (llama/20480)
lhez [Sat, 14 Mar 2026 05:18:52 +0000 (22:18 -0700)]
opencl: fix l2_norm (llama/20480)

5 weeks agograph : remove redundant GDN state transposes (llama/20443)
Georgi Gerganov [Fri, 13 Mar 2026 20:12:54 +0000 (22:12 +0200)]
graph : remove redundant GDN state transposes (llama/20443)

* ggml : transpose fused GDN state access for coalesced memory reads (llama/20436)

The fused Gated Delta Net kernel accessed the [S_v, S_v] state matrix
column-wise on row-major storage, causing strided reads (stride S_v =
128 floats = 512 bytes) that waste GPU cache bandwidth. This produced a
39% regression on Qwen3.5-9B (Metal, M4 Max) compared to the unfused
path.

Transpose the state indexing so threads read contiguously:
- Metal: s_ptr[is*S_v] -> s_ptr[is] (stride 1 vs S_v)
- CUDA:  curr_state[i*S_v+col] -> curr_state[col*S_v+i] (coalesced)
- CPU:   restructured loops for row-wise transposed access

Also add --fused-gdn [on|off|auto] CLI flag (mirrors --flash-attn) so
users can control fused GDN independently of auto-detection.

All GATED_DELTA_NET backend-ops tests pass.

Co-Authored-By: Claude Opus 4.6 <redacted>
* ggml : use SIMD dot products in CPU GDN kernel, couple AR/chunked fused flags

- Replace scalar inner loops with ggml_vec_dot_f32 for SIMD-optimized
  dot products in the CPU fused GDN kernel (delta and attention output)
- Couple fused_gdn_ar and fused_gdn_ch flags in auto-detection: if one
  path lacks device support, disable both to prevent state layout mismatch
  between transposed (fused) and non-transposed (unfused) formats

Co-Authored-By: Claude Opus 4.6 <redacted>
* llama : rever fgdn argument changes

* graph : remove GDN state transposes

* vulkan : adapt

* cuda : remove obsolete smem code

---------

Co-authored-by: Paul Flynn <redacted>
Co-authored-by: Claude Opus 4.6 <redacted>
Co-authored-by: Oliver Simons <redacted>
5 weeks agoggml-cpu: add RVV vec dot kernels for quantization types (llama/18859)
rehan-10xengineer [Fri, 13 Mar 2026 15:36:04 +0000 (20:36 +0500)]
ggml-cpu: add RVV vec dot kernels for quantization types (llama/18859)

* ggml-cpu: add rvv quantize_row_q8_K kernel

Co-authored-by: Rehan Qasim <redacted>
* ggml-cpu: add rvv vec_dot for iq4_nl, mxfp4, iq2_xxs

Co-authored-by: Rehan Qasim <redacted>
* ggml-cpu: add rvv vec_dot for iq4_xs, refactor

* ggml-cpu: remove ifunc for rvv vec dot

* ggml-cpu: add vec_dot for iq2_xs, iq3_xxs

Co-authored-by: Rehan Qasim <redacted>
* ggml-cpu: refactor quants.c

---------

Co-authored-by: taimur-10x <redacted>
Co-authored-by: Rehan Qasim <redacted>
Co-authored-by: Rehan Qasim <redacted>
5 weeks agoggml : fix typo gmml (llama/20512)
Adrien Gallouët [Fri, 13 Mar 2026 13:36:13 +0000 (14:36 +0100)]
ggml : fix typo gmml (llama/20512)

Signed-off-by: Adrien Gallouët <redacted>
5 weeks agometal : fix l2 norm scale (llama/20493)
Georgi Gerganov [Fri, 13 Mar 2026 09:43:20 +0000 (11:43 +0200)]
metal : fix l2 norm scale (llama/20493)

5 weeks agollama : disable graph reuse with pipeline parallelism (llama/20463)
Georgi Gerganov [Thu, 12 Mar 2026 19:04:13 +0000 (21:04 +0200)]
llama : disable graph reuse with pipeline parallelism (llama/20463)

5 weeks agovulkan: add GATED_DELTA_NET op support (llama/20334)
ProgenyAlpha [Thu, 12 Mar 2026 10:32:04 +0000 (06:32 -0400)]
vulkan: add GATED_DELTA_NET op support (llama/20334)

* vulkan: add GATED_DELTA_NET op support

Implements the fused gated delta net recurrence as a Vulkan compute
shader with full support for scalar gate, KDA vector gate, GQA
broadcast, multi-token sequences, and permuted (non-contiguous) q/k
inputs. Specialization constants select head size (32/64/128) and
KDA mode at pipeline creation time.

Passes all 13 test-backend-ops cases on AMD Radeon 890M (RADV GFX1150).

Co-Authored-By: Claude Opus 4.6 <redacted>
* vulkan: optimize GATED_DELTA_NET shader (Phase 1)

- vec4 dot products on all inner loops (dp4 hardware intrinsic)
- Cache exp(g) in shared memory for KDA path, eliminating ~32K
  redundant global reads and ~16K redundant exp() calls per token
- vec4 fused decay + rank-1 update (3 vec4 ops vs 12 scalar ops)
- Add perf benchmark cases for GATED_DELTA_NET to test-backend-ops

KDA TG: +5.4% throughput. Non-KDA: no regressions.
13/13 test-backend-ops passing on AMD Radeon 890M (RADV GFX1150).

Co-Authored-By: Claude Opus 4.6 <redacted>
* vulkan: address review feedback for GATED_DELTA_NET

Pipeline array refactor [3][2], A_TYPE/D_TYPE/FLOAT_TYPE shader macros,
scale in push constants, supports_op fix, dispatch restructuring.

Co-Authored-By: Claude Opus 4.6 <redacted>
* vulkan: use FLOAT_TYPE for buffer/shared declarations, align formatting

Co-Authored-By: Claude Opus 4.6 <redacted>
* vulkan: add explicit FLOAT_TYPE casts for buffer loads

Wrap data_q, data_k, and data_g buffer reads with FLOAT_TYPE() casts
to ensure correct behavior across all Vulkan configurations.

Co-Authored-By: Claude Opus 4.6 <redacted>
* vulkan: fix Q/K broadcast for interleaved head layout

Adapt to the interleaved broadcast convention from #20340:
head_id / rq1 → head_id % neq1

Co-Authored-By: Claude Opus 4.6 <redacted>
---------

Co-authored-by: Progeny Alpha <redacted>
Co-authored-by: Claude Opus 4.6 <redacted>
5 weeks agovulkan: fix SSM_CONV PP scaling with large ubatch sizes (llama/20379)
ProgenyAlpha [Thu, 12 Mar 2026 09:03:18 +0000 (05:03 -0400)]
vulkan: fix SSM_CONV PP scaling with large ubatch sizes (llama/20379)

* vulkan: optimize SSM_CONV workgroup dispatch for large ubatch

Tile tokens into 2D workgroups (32x16) to reduce workgroup launch
overhead at large ubatch sizes. Add vec4 fast path for nc=4 (common
d_conv size). Fixes PP performance degradation with ubatch > 512.

Ref: ggml-org/llama.cpp#18725

Co-Authored-By: Claude Opus 4.6 <redacted>
* vulkan: remove unused shared memory declaration in SSM_CONV

Co-Authored-By: Claude Opus 4.6 <redacted>
---------

Co-authored-by: Progeny Alpha <redacted>
Co-authored-by: Claude Opus 4.6 <redacted>
5 weeks agosync : ggml
Georgi Gerganov [Mon, 16 Mar 2026 05:13:14 +0000 (07:13 +0200)]
sync : ggml

5 weeks agometal : avoid divisions in bin kernel (llama/20426)
Georgi Gerganov [Mon, 16 Mar 2026 05:12:50 +0000 (07:12 +0200)]
metal : avoid divisions in bin kernel (llama/20426)

5 weeks agosync : ggml
Georgi Gerganov [Mon, 16 Mar 2026 05:12:37 +0000 (07:12 +0200)]
sync : ggml

5 weeks agovulkan: fix l2_norm epsilon handling (llama/20350)
Jeff Bolz [Thu, 12 Mar 2026 05:39:41 +0000 (00:39 -0500)]
vulkan: fix l2_norm epsilon handling (llama/20350)

5 weeks agovulkan: fix OOB check in flash_attn_mask_opt (llama/20296)
Jeff Bolz [Thu, 12 Mar 2026 05:35:49 +0000 (00:35 -0500)]
vulkan: fix OOB check in flash_attn_mask_opt (llama/20296)

5 weeks agovulkan: Fix ErrorOutOfHostMemory on Intel GPU when loading large models with --no...
Masato Nakasaka [Thu, 12 Mar 2026 05:30:16 +0000 (22:30 -0700)]
vulkan: Fix ErrorOutOfHostMemory on Intel GPU when loading large models with --no-mmap (llama/20059)

* Changed to reuse command buffers to fix crashing on Intel GPU

* Removed unused parameter

* Fixed compile error and minor mistake

* Fix logging

* Changing to use usage flag per command buffer

* fixed style

* added buffer reset

* Removed cmd_buffer_idx for reuse consistency

* Fixed style

5 weeks agoopencl: use larger workgroup size for get_rows (llama/20316)
lhez [Thu, 12 Mar 2026 05:03:27 +0000 (22:03 -0700)]
opencl: use larger workgroup size for get_rows (llama/20316)

5 weeks agoopencl: add cumsum op (llama/18981)
shaofeiqi [Thu, 12 Mar 2026 05:03:07 +0000 (22:03 -0700)]
opencl: add cumsum op (llama/18981)

* OpenCL: add CUMSUM op support

* remove unused argument

* opencl: refactor cumsum

* opencl: refactor

* opencl: refactor tmp buffer

* opencl: adjust max number of subgroups

* opencl: fix whitespace

* opencl: fix global size when cumsum the tmp buffer

---------

Co-authored-by: Li He <redacted>
5 weeks agohip: compile debug builds with -O2 on hip to avoid a compiler bug (llama/20392)
uvos [Thu, 12 Mar 2026 02:37:10 +0000 (03:37 +0100)]
hip: compile debug builds with -O2 on hip to avoid a compiler bug (llama/20392)

5 weeks agoggml-webgpu: Add supports for `GGML_OP_REPEAT` (llama/20230)
Masashi Yoshimura [Wed, 11 Mar 2026 21:40:36 +0000 (06:40 +0900)]
ggml-webgpu: Add supports for `GGML_OP_REPEAT` (llama/20230)

* Add GGML_OP_REPEAT to webgpu backend.

* Add i16 support for GGML_OP_REPEAT.

5 weeks agollama : enable chunked fused GDN path (llama/20340)
Georgi Gerganov [Wed, 11 Mar 2026 20:46:40 +0000 (22:46 +0200)]
llama : enable chunked fused GDN path (llama/20340)

* llama : enable chunked fused GDN path

* models : avoid Q and K repeats when using fused GDA

* cont : fix comment

Co-authored-by: Aman Gupta <redacted>
* cont : fix the fix

Co-authored-by: Aman Gupta <redacted>
* cont : fix

* metal : add GDN kernel (llama/20361)

* metal : add Metal backend for GGML_OP_GATED_DELTA_NET

Add a fused Metal kernel for the gated delta net recurrence op
(#19504), enabling GPU-accelerated inference for DeltaNet-based
models (Qwen3.5, etc.) on Apple Silicon.

Supports both GDA (scalar gate) and KDA (per-row gate) modes
with head_size 64 and 128. Unsupported configurations (head_size
32, non-contiguous tensors) gracefully fall back to CPU.

Performance: Qwen3.5-0.8B Q4_K_M on M4 Max
  tg128: 170 -> 213 t/s (+25%)

Co-Authored-By: Claude Opus 4.6 <redacted>
* metal : validate contiguity of all input tensors in supports_op

Co-Authored-By: Claude Opus 4.6 <redacted>
* metal : add algorithm equivalence comment for GDA decay path

Co-Authored-By: Claude Opus 4.6 <redacted>
* cont : unslop + optimize

* cont : clean-up

---------

Co-authored-by: Paul Flynn <redacted>
Co-authored-by: Claude Opus 4.6 <redacted>
* CUDA: AR gated delta net improvements (llama/20391)

* Add FastDiv to gated_delta_net_cuda

* Shard columns across warps

This reduces register pressure (avoids spill for S_v = 128) and gives
the warp-scheduler more CTAs to schedule (thus hiding data-access
latencies).

* Remove unneded include in gated_delta_net.cu

* Improve comments

* Apply code-formating

* Make sharding HIP-compatible

1. Use ggml_cuda_get_physical_warp_size() to determine warp size flexibly
2. Add test with partial warp to test sum reduction on CUDA

* Remove fastdiv_s64, as we can treat neqk1 and rq3 as uint32_t

* Rename variables

* Enable GDN also for prefill, move TODO for chunked_GDN

* Actually remove the TODO from 206890897546bd16602c3b79394fd5ea09ef199f

* Get warp size at runtime

warp_size is not known at compile time in hip host code.

* Don't expose ggml_cuda_get_physical_warp_size on host

---------

Co-authored-by: uvos <redacted>
* llama : refactor llm_build_delta_net_base API

---------

Co-authored-by: Aman Gupta <redacted>
Co-authored-by: Paul Flynn <redacted>
Co-authored-by: Claude Opus 4.6 <redacted>
Co-authored-by: Oliver Simons <redacted>
Co-authored-by: uvos <redacted>
5 weeks agoggml : add NVFP4 quantization type support (llama/19769)
Richard Davison [Wed, 11 Mar 2026 20:02:54 +0000 (21:02 +0100)]
ggml : add NVFP4 quantization type support (llama/19769)

* WIP: add NVFP4 quantization support

* tests

* improve NVFP4 dot product implementation performance and fix bad super call

* typo

* Use nvfp4 kvalues

* vulkan : fix NVFP4 shader compilation by including kvalues_mxfp4 lookup table

* vulcal and perf fixes

* wip

* Fix metal

* fix vulcan

* Rename threshold & fix wrong scale

* Fix MOE

* Shelf backend implementations (CUDA, Metal, Vulkan, arch-specific SIMD)

Remove NVFP4 support from GPU backends and architecture-specific
optimized dot products. These should be added in separate PRs so
backend specialists can review them independently.

Reverted files:
- ggml-cuda: common.cuh, convert.cu, mmq.cu/cuh, mmvq.cu, vecdotq.cuh,
  quantize.cu/cuh, mma.cuh, ggml-cuda.cu, fattn-tile.cuh
- ggml-metal: ggml-metal.metal, ggml-metal-device.cpp, ggml-metal-impl.h,
  ggml-metal-ops.cpp
- ggml-vulkan: ggml-vulkan.cpp, all vulkan-shaders/*
- ggml-cpu arch: arm/quants.c, x86/quants.c, powerpc/quants.c, s390/quants.c

Core NVFP4 support (type definition, CPU fallback dot product,
quantization, dequantization, conversion) is retained.

* Fix arch-fallback.h: add NVFP4 generic fallback for all platforms

After shelving backend-specific SIMD implementations, the generic
CPU dot product needs to be aliased on ARM, x86, PowerPC, and s390
platforms that previously relied on arch-specific versions.

* quantize: add NVFP4 as a quantization type option

* Fix ggml_fp32_to_ue4m3: handle subnormal values

Previously, values with ue4m3_exp <= 0 were clamped to 0, causing
all small scales to underflow. This made NVFP4 quantization via
llama-quantize produce garbage (PPL = 5.8M) since typical transformer
weights have amax/6.0 in the range 0.001-0.01, which falls in the
UE4M3 subnormal range.

Now subnormals are properly encoded as man * 2^-9 (exp=0, man=1..7),
matching the decode path in ggml_ue4m3_to_fp32.

Result: NVFP4 requantization now produces PPL = 15.25 (vs F16 = 14.33),
comparable to Q4_1 (PPL = 15.81) at slightly lower BPW (4.70 vs 5.15).

* Restore ARM NEON NVFP4 dot product implementation

Restores the optimized ggml_vec_dot_nvfp4_q8_0 for ARM NEON using
vqtbl1q_s8 lookup and ggml_vdotq_s32 dot products.

tg128 performance: 4.37 t/s (generic) -> 13.66 t/s (NEON) = 3.1x speedup

* Optimize ARM NEON NVFP4 dot product: LUT + vpaddq + vfmaq

- Add ue4m3_scale_lut[128] to ggml-common.h replacing branch-heavy
  ggml_ue4m3_to_fp32() in the hot loop
- Use vpaddq_s32 for pairwise int32 reduction instead of vaddvq_s32
- Accumulate with vfmaq_f32 into float32x4_t vector accumulators

tg128: 8.1 -> 31.0 t/s (3.8x speedup, 77% of Q4_1 speed)

* ARM NEON NVFP4: rearrange q8 to match nibble layout

Alternative approach: rearrange q8 data to match the NVFP4 lo/hi
nibble layout instead of rearranging the looked-up NVFP4 values.
Eliminates vcombine_s8(vget_low, vget_low) shuffles.

Performance is equivalent (~18.5 t/s) - the bottleneck is the 2x
block overhead from QK=16 vs QK=32, not the shuffle instructions.

* CPU only backend 64 super-block layout

* cleanup

* Remove unused LUT

* int

* exclude NVFP4 from unsupported ops in metal build

* remove quantization for now

* store scales as native UE4M3, preserve original model bits when possible

* Update convert_hf_to_gguf.py

Co-authored-by: Sigbjørn Skjæret <redacted>
* correct comment

* format

* reduce duplication and cleanup

* Address comments

* move detection to prepare_tensors

* Use math instead of const

* Move

* fix comment

* Shelf quantize tests

* Rebase and move check

* cleanup

* lint

* Update gguf-py/gguf/scripts/gguf_convert_endian.py

Co-authored-by: Sigbjørn Skjæret <redacted>
* Use fallback quant config

* Simplify

Co-authored-by: Sigbjørn Skjæret <redacted>
* organize

* Refactor

* Update convert_hf_to_gguf.py

Co-authored-by: Sigbjørn Skjæret <redacted>
* Update convert_hf_to_gguf.py

Co-authored-by: Sigbjørn Skjæret <redacted>
* Update convert_hf_to_gguf.py

Co-authored-by: Sigbjørn Skjæret <redacted>
* add quantize_nvfp4 (required for test_quants.py)

* add quantize_nvfp4 (required for test_quants.py)

* add quantize_nvfp4 (required for test_quants.py)

* fix return type

---------

Co-authored-by: Sigbjørn Skjæret <redacted>
5 weeks agollama : add support for Nemotron 3 Super (llama/20411)
Daniel Bevenius [Wed, 11 Mar 2026 18:27:53 +0000 (19:27 +0100)]
llama : add support for Nemotron 3 Super (llama/20411)

* llama : add support for Nemotron 3 Super

This commit adds support for the Nemotron 3 Super model (120B.A12B)
enabling this model to be converted to GGUF format and run in llama.cpp.

Co-authored-by: Georgi Gerganov <redacted>
Co-authored-by: Matt Clayton <redacted>
5 weeks agometal : fix capture_compute counter logic (llama/20410)
Georgi Gerganov [Wed, 11 Mar 2026 16:38:22 +0000 (18:38 +0200)]
metal : fix capture_compute counter logic (llama/20410)

5 weeks agometal : fix q5_k mul_mv register spill (llama/20399)
Georgi Gerganov [Wed, 11 Mar 2026 14:25:27 +0000 (16:25 +0200)]
metal : fix q5_k mul_mv register spill (llama/20399)

5 weeks agometal : add env var to trigger graph capture (llama/20398)
Georgi Gerganov [Wed, 11 Mar 2026 14:25:10 +0000 (16:25 +0200)]
metal : add env var to trigger graph capture (llama/20398)

5 weeks agoggml-cuda: gdn use shared mem for HIP (llama/20366)
uvos [Wed, 11 Mar 2026 05:06:19 +0000 (06:06 +0100)]
ggml-cuda: gdn use shared mem for HIP (llama/20366)

Suggested-by: Aman Gupta <redacted>
5 weeks agocuda/hip: fix loop unrolling in ssm-conv (llama/20369)
uvos [Wed, 11 Mar 2026 05:04:32 +0000 (06:04 +0100)]
cuda/hip: fix loop unrolling in ssm-conv (llama/20369)

5 weeks agofix op rope, add rope_back (llama/20293)
Neo Zhang [Wed, 11 Mar 2026 01:53:34 +0000 (09:53 +0800)]
fix op rope, add rope_back (llama/20293)

5 weeks agofix for failed UT case: ACC, L2_NORM, UPSCALE, fused_glu, unary (llama/20283)
Neo Zhang [Wed, 11 Mar 2026 01:53:05 +0000 (09:53 +0800)]
fix for failed UT case: ACC, L2_NORM, UPSCALE, fused_glu, unary (llama/20283)

5 weeks agoggml : bump RPC version (llama/20330)
Georgi Gerganov [Tue, 10 Mar 2026 19:36:57 +0000 (21:36 +0200)]
ggml : bump RPC version (llama/20330)

5 weeks agoggml webgpu: faster normal quant and some k-quant matrix operations, better shader...
Reese Levine [Tue, 10 Mar 2026 16:14:27 +0000 (09:14 -0700)]
ggml webgpu: faster normal quant and some k-quant matrix operations, better shader parameter handling (llama/20173)

* K quant speedup (llama/20)

* Basic JIT compilation for mul_mat, get_rows, and scale (llama/17)

* scale jit working

* preliminary working jit for getrows and mulmat, needs refining

* simplified mul_mat preprocessing switch statement

* get_rows fixes, mul_mat refinement

* formatted + last edits

* removed some extraneous prints

* fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish

* small fix

* some changes, working

* get_rows and mul_mat jit fixed and working

* Update formatting

* formatting

* Add header

---------

Co-authored-by: Neha Abbas <redacted>
Co-authored-by: Reese Levine <redacted>
* Start work on all-encompassing shader library

* refactor argmax, set_rows

* Refactor all but flashattention, mat mul

* no gibberish, all k quants added, merged

* vec memory fix

* q6_k matching metal on my machine, tests passing

* Set tile size for q6_k separately

* Separate out fast shaders

---------

Co-authored-by: neha-ha <redacted>
* Move towards writeBuffer for params

* Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups

* Remove extra file

* Formatting

---------

Co-authored-by: neha-ha <redacted>
5 weeks agokleidiai : support for concurrent sme and neon kernel execution (llama/20070)
Charles Xu [Tue, 10 Mar 2026 07:25:25 +0000 (08:25 +0100)]
kleidiai : support for concurrent sme and neon kernel execution (llama/20070)

5 weeks agoggml-cpu: add RVV repack GEMM and GEMV for quantization types (llama/19121)
Taimur Ahmad [Tue, 10 Mar 2026 06:49:52 +0000 (11:49 +0500)]
ggml-cpu: add RVV repack GEMM and GEMV for quantization types (llama/19121)

* ggml-cpu: add rvv ggml_quantize_mat_4x8 for q8_0

Co-authored-by: Rehan Qasim <redacted>
* ggml-cpu: add rvv repacking for iq4_nl

* ggml-cpu: add generic impl for iq4_nl gemm/gemv

* ggml-cpu: add rvv repacking for q8_0

* ggml-cpu: refactor; add rvv repacking for q4_0, q4_K

* ggml-cpu: refactor; add rvv repacking for q2_K

Co-authored-by: Rehan Qasim <redacted>
* ggml-cpu: refactor rvv repack

---------

Co-authored-by: Rehan Qasim <redacted>
5 weeks agometal: handle command buffer failures gracefully in synchronize (llama/20306)
Julian Pscheid [Tue, 10 Mar 2026 06:32:24 +0000 (23:32 -0700)]
metal: handle command buffer failures gracefully in synchronize (llama/20306)

Replace GGML_ABORT("fatal error") in ggml_metal_synchronize() with
error flag + return. This aligns synchronize error handling with
graph_compute, which already returns GGML_STATUS_FAILED for the same
condition.

When a command buffer fails (e.g., iOS GPU access revocation during
backgrounding, macOS eGPU disconnect, OOM), the backend enters an
error state instead of killing the host process. Subsequent
graph_compute calls return GGML_STATUS_FAILED immediately. Recovery
requires recreating the backend.

Failed extra command buffers are properly released on the error path
to avoid Metal object leaks.

5 weeks agometal : extend mul_mv_ext to BF16, Q2_K, Q3_K (llama/20250)
Paul Flynn [Mon, 9 Mar 2026 14:48:12 +0000 (10:48 -0400)]
metal : extend mul_mv_ext to BF16, Q2_K, Q3_K (llama/20250)

Enable mul_mv_ext small-batch kernels (BS 2-8) for BF16, Q2_K,
and Q3_K quantization types. These types previously fell through
to the slower single-row mul_mv path.

BF16 uses the float4 dequantize path (like F16). Q2_K and Q3_K
use the float4x4 K-quant path (like Q4_K/Q5_K/Q6_K).

Co-authored-by: Claude Opus 4.6 <redacted>
5 weeks agometal : add upscale (llama/20284)
Georgi Gerganov [Mon, 9 Mar 2026 14:45:11 +0000 (16:45 +0200)]
metal : add upscale (llama/20284)

5 weeks agoggml-cuda: disable gdn for musa (llama/20278)
Aman Gupta [Mon, 9 Mar 2026 08:15:36 +0000 (16:15 +0800)]
ggml-cuda: disable gdn for musa (llama/20278)

5 weeks agoggml-vulkan: add SGN operator, auto-generate Vulkan.csv and ops.md (llama/20219)
Bertay Eren [Mon, 9 Mar 2026 06:24:16 +0000 (09:24 +0300)]
ggml-vulkan: add SGN operator, auto-generate Vulkan.csv and ops.md (llama/20219)

5 weeks agovulkan: skip zero size tensors in backend copies (llama/20233)
Ruben Ortlam [Mon, 9 Mar 2026 06:23:45 +0000 (07:23 +0100)]
vulkan: skip zero size tensors in backend copies (llama/20233)

5 weeks agocuda : display total and free VRAM capacity during device initialization (llama/20185)
Michael Huang [Mon, 9 Mar 2026 04:45:43 +0000 (21:45 -0700)]
cuda : display total and free VRAM capacity during device initialization (llama/20185)

5 weeks agoggml-vulkan: Add ELU op support (llama/20183)
GiantPrince [Sun, 8 Mar 2026 11:38:17 +0000 (07:38 -0400)]
ggml-vulkan: Add ELU op support (llama/20183)

* ggml-Vulkan: add ELU support

* ggml-Vulkan: remove extra spaces and variables

* ggml-Vulkan: fix format issue

* ggml-Vulkan: fix format issue

* fix whitespace issue

* Update Vulkan.csv and ops.md

5 weeks agovulkan: Fix data races in coopmat1 mul_mat(_id) (llama/20084)
Jeff Bolz [Sun, 8 Mar 2026 11:33:48 +0000 (06:33 -0500)]
vulkan: Fix data races in coopmat1 mul_mat(_id) (llama/20084)

* vulkan: Fix data races in coopmat1 mul_mat(_id)

Add barriers between coopmat store and regular loads. We sort of got away with
this because it was the same subgroup accessing the values, but it's still a
race and may not work.

* switch to subgroup control barriers

5 weeks agosupprt Flash Attention for fp32/fp16/Q4/Q5/Q8 (llama/20190)
Neo Zhang [Sun, 8 Mar 2026 04:00:07 +0000 (12:00 +0800)]
supprt Flash Attention for fp32/fp16/Q4/Q5/Q8 (llama/20190)

* support flash-attention for fp32/fp16/Q4/Q5/Q8

* rm warining

* update for JIT

5 weeks agoggml: add GATED_DELTA_NET op (llama/19504)
Aman Gupta [Sat, 7 Mar 2026 07:41:10 +0000 (15:41 +0800)]
ggml: add GATED_DELTA_NET op (llama/19504)

* ggml: add GATED_DELTA_NET op

* remove the transpose

* add KDA

* add qwen35 dense

* llama : check for fused gated delta net backend support

---------

Co-authored-by: Georgi Gerganov <redacted>
5 weeks agoopencl: add l2_norm (llama/20160)
lhez [Sat, 7 Mar 2026 02:03:05 +0000 (18:03 -0800)]
opencl: add l2_norm (llama/20160)

5 weeks agoquants : Add memsets and other fixes for IQ quants (llama/19861)
Bartowski [Fri, 6 Mar 2026 21:06:56 +0000 (16:06 -0500)]
quants : Add memsets and other fixes for IQ quants (llama/19861)

* Add memsets and other fixes for IQ quants

* Make memset unconditional, change Laux back to L

* Move another memset

5 weeks agohexagon: add f32 ssm_conv op (llama/20122)
Todor Boinovski [Fri, 6 Mar 2026 17:59:26 +0000 (09:59 -0800)]
hexagon: add f32 ssm_conv op (llama/20122)

* hexagon: add ssm_conv op

* hexagon: hvx kernel is functional

* hexagon: improvements to ssm-conv hvx kernel

* hexagon: added dma to ssm-conv hvx kernel

* hexagon: ssm-conv dynamically compute gather scratchpad

* hex-ssm-conv: add local context and fix various issues (spad indexing, etc)

---------

Co-authored-by: Max Krasnyansky <redacted>
5 weeks agocpu: skip redudant ROPE cache updates (llama/20149)
Max Krasnyansky [Fri, 6 Mar 2026 16:32:40 +0000 (08:32 -0800)]
cpu: skip redudant ROPE cache updates (llama/20149)

5 weeks agoggml-cuda: add mem check for fusion (llama/19916)
Aman Gupta [Fri, 6 Mar 2026 16:05:43 +0000 (00:05 +0800)]
ggml-cuda: add mem check for fusion (llama/19916)

* ggml-cuda: add mem check for fusion

* Replace NaNs with -FLT_MAX

* fix typo

Co-authored-by: Johannes Gäßler <redacted>
---------

Co-authored-by: Johannes Gäßler <redacted>
5 weeks agoggml: update comments for backends which have no memory to report (llama/20157)
Aaron Teo [Fri, 6 Mar 2026 15:24:38 +0000 (23:24 +0800)]
ggml: update comments for backends which have no memory to report (llama/20157)

Signed-off-by: Aaron Teo <redacted>
5 weeks agoggml-cpu: Fix gcc 15 ICE on ppc64le (ggml/20083) (llama/20130)
shalinib-ibm [Fri, 6 Mar 2026 15:22:39 +0000 (20:52 +0530)]
ggml-cpu: Fix gcc 15 ICE on ppc64le (ggml/20083) (llama/20130)

This patch addresses an Internal Compiler Error (Segmentation fault)
observed with gcc 15 by replacing the intrinsic + cast by doing
a cat on the data first and then calling the intrinsic. This bypasses the
buggy compiler path while maintaining identical instruction selection.

Performance Verification:
Assembly analysis on RHEL 9 (GCC 15.1.1) confirms that both the original
code and this fix generate the identical Power10 prefixed load instruction:
    `plxv 40, 2(14)`

This ensures zero performance regression while unblocking builds on
newer toolchains.

Reproduced on:
- Alpine Linux + GCC 15.2.0-r2
- RHEL 9  + GCC 15.1.1 (gcc-toolset-15)

Signed-off-by: Shalini Salomi Bodapati <redacted>
5 weeks agoCUDA: use shared mem for ssm_conv (llama/20128)
Aman Gupta [Fri, 6 Mar 2026 15:09:59 +0000 (23:09 +0800)]
CUDA: use shared mem for ssm_conv (llama/20128)

* CUDA: use shared mem for ssm_conv

* fuse silu + ssm_conv

* fuse unary + mul

* enable for fp16

* formatting

Co-authored-by: Johannes Gäßler <redacted>
---------

Co-authored-by: Johannes Gäßler <redacted>
5 weeks agoggml-cpu: fix data race for debug asserts (llama/20148)
Johannes Gäßler [Fri, 6 Mar 2026 08:12:49 +0000 (09:12 +0100)]
ggml-cpu: fix data race for debug asserts (llama/20148)

5 weeks agoopencl: add neg, exp and diag (llama/20127)
lhez [Fri, 6 Mar 2026 05:16:39 +0000 (21:16 -0800)]
opencl: add neg, exp and diag (llama/20127)

* opencl: add `neg`

* opencl: add `exp`

* opencl: add `diag`

5 weeks agohexagon: add fp16 support for binary ops: add,sub,mul,div (llama/20139)
YardenTal44 [Fri, 6 Mar 2026 02:29:13 +0000 (04:29 +0200)]
hexagon: add fp16 support for binary ops: add,sub,mul,div (llama/20139)

* hexagon: add fp16 support for binary ops: add,sub,mul,div

* hexagon: fix test-backend-ops failures for fp16 binary ops on older arches (<v79)

* hexagon: decide on n_threads (aka n_jobs) early to avoid overallocating scratchpad

* snapdragon: fix readme link

---------

Co-authored-by: Max Krasnyansky <redacted>
5 weeks agoCUDA: Improve performance via less synchronizations between token (llama/17795)
Andreas Kieslinger [Thu, 5 Mar 2026 11:53:21 +0000 (12:53 +0100)]
CUDA: Improve performance via less synchronizations between token (llama/17795)

* Adds CPU-to-CUDA copy capability to
ggml_backend_cuda_cpy_tensor_async()

* Adds function to relax sync requirements between input copies on
supported backends (CUDA for now)

* Exchanges synchronous copy with async copy function.

* Adds macro guards to allow compilation in non-CUDA builds

* Reworked backend detection in ggml-backend.cpp to avoid linking
conflicts

* Relax requirement of checks in async CUDA copies from backend and buffer type to just buffer type, to avoid linking issues

* Minor cleanup

* Makes opt-in to relax use of explicit syncs more general. Backends like
vulkan which require a synchronization between HtoD copies and graph
execution could also adopt this change now.

* Reintroduces stricter check for CPU->CUDA backend async copy via
GGML_DEVICE_TYPE_CPU.

* Corrects initialization of ggml_backend_sync_mode in
ggml_backend_sched_split initialization

* Simplifies synchronizations to adhere to `saaasg` pattern.

* Apply suggestion from @ggerganov (src->buffer to buf_src)

Co-authored-by: Georgi Gerganov <redacted>
* Apply suggestion from @ggerganov (src->buffer to buf_src) v2

Co-authored-by: Georgi Gerganov <redacted>
---------

Co-authored-by: Georgi Gerganov <redacted>
5 weeks agochore : correct typos [no ci] (llama/20041)
Marcel Petrick [Thu, 5 Mar 2026 07:50:21 +0000 (08:50 +0100)]
chore : correct typos [no ci] (llama/20041)

* fix(docs): correct typos found during code review

Non-functional changes only:
- Fixed minor spelling mistakes in comments
- Corrected typos in user-facing strings
- No variables, logic, or functional code was modified.

Signed-off-by: Marcel Petrick <redacted>
* Update docs/backend/CANN.md

Co-authored-by: Aaron Teo <redacted>
* Revert "Auxiliary commit to revert individual files from 846d1c301281178efbc6ce6060ad34c1ebe45af8"

This reverts commit 02fcf0c7db661d5ff3eff96b2b2db9fdb7213256.

* Update tests/test-backend-ops.cpp

Co-authored-by: Sigbjørn Skjæret <redacted>
* Update tests/test-backend-ops.cpp

Co-authored-by: Sigbjørn Skjæret <redacted>
---------

Signed-off-by: Marcel Petrick <redacted>
Co-authored-by: Aaron Teo <redacted>
Co-authored-by: Sigbjørn Skjæret <redacted>
5 weeks agohexagon: Flash Attention optimizations (dma, mpyacc, multi-row) and MatMul updates...
Max Krasnyansky [Thu, 5 Mar 2026 05:55:29 +0000 (21:55 -0800)]
hexagon: Flash Attention optimizations (dma, mpyacc, multi-row) and MatMul updates (llama/20118)

* ggml-hexagon: enhance hvx_dot_f16_f16_aa_rx4 for improved performance by expanding vector handling and optimizing accumulation

# Conflicts:
# ggml/src/ggml-hexagon/htp/flash-attn-ops.c

* ggml-hexagon: optimize hvx_dot_f16_f16_aa_rx4 and enhance hvx_vec_reduce_sum_f32x4 for improved performance and reduced complexity

* ggml-hexagon: add hvx_dot_f16_f16_aa_rx32 for enhanced vector processing in flash attention

# Conflicts:
# ggml/src/ggml-hexagon/htp/flash-attn-ops.c

* optimize hvx_dot_f16_f16_aa_rx4 and hvx_dot_f16_f16_aa_rx32 by removing unused scale parameter and improving vector accumulation

# Conflicts:
# ggml/src/ggml-hexagon/htp/flash-attn-ops.c

* ggml-hexagon: refactor hvx_dot_f16_f16_aa_rx4 for improved readability and return HVX_Vector for better integration

# Conflicts:
# ggml/src/ggml-hexagon/htp/flash-attn-ops.c

* ggml-hexagon: initialize sums variable in hvx_dot_f16_f16_aa_rx32 for clarity

* ggml-hexagon: fix compiling error

* fix hvx_dot_f16_f16_aa_rx4 to handle leftover elements correctly using masking

* refactor hvx_dot_f16_f16_aa_rx4 to accept vector and leftover element counts as parameters for improved clarity and flexibility

* wip

* fa: instrumentation and dma reordering

* hex-fa: use block-size 64 to improve DMA pipelining

* hex-fa: optimize vec-dot for v79 and above

* hex-fa: use block size 64

* hex-fa: avoid scalar fp32->fp16 conversions

* hex-fa: simplify dot_f16 functions using optimized vec_mpyacc

* hex-fa: rewrite mad_f32_f16 using hvx_vec_mpyacc

* hex-mm: use mpyacc in matmul dot functions

---------

Co-authored-by: chraac <redacted>
5 weeks agoopencl: add `SET`, support i32 for `CPY`, minor refactor for cpy (llama/20101)
lhez [Thu, 5 Mar 2026 05:32:26 +0000 (21:32 -0800)]
opencl: add `SET`, support i32 for `CPY`, minor refactor for cpy (llama/20101)

5 weeks agoFix wait logic for inflight jobs (llama/20096)
Nikhil Jain [Wed, 4 Mar 2026 19:54:55 +0000 (11:54 -0800)]
Fix wait logic for inflight jobs (llama/20096)

* Enable tmate debugging for investigating thread safety issue

* Refactor wait and submit to operate on vector<wgpu::FutureWaitInfo>, and fix wait to delete only the future that is completed.

* Cleanup

* Remove clear change and run clang-format

* Cleanup

5 weeks agoAdd concat op to webgpu. (llama/20068)
Masashi Yoshimura [Wed, 4 Mar 2026 19:19:00 +0000 (04:19 +0900)]
Add concat op to webgpu. (llama/20068)

5 weeks agoggml: fix ggml_is_contiguous_n for ne == 1 (llama/20092)
Johannes Gäßler [Wed, 4 Mar 2026 11:04:31 +0000 (12:04 +0100)]
ggml: fix ggml_is_contiguous_n for ne == 1 (llama/20092)

5 weeks agoggml : use a simple std::thread in AMX without OpenMP (llama/20074)
Adrien Gallouët [Wed, 4 Mar 2026 10:57:09 +0000 (11:57 +0100)]
ggml : use a simple std::thread in AMX without OpenMP (llama/20074)

Disabling OpenMP generally provides better inference performance (at
least in my testing) but the loading becomes slightly slower.

Benchmark results for `convert_B_packed_format()`:

Before this commit:

         N      K |  No OpenMP     OpenMP |    Diff |  Speedup
    ------------------------------------------------------------
       512   2880 |    640.9us    263.5us |  -58.9% |    0.41x
      2880   4096 |     2.55ms    261.7us |  -89.8% |    0.10x
    201088   2880 |   256.44ms    21.61ms |  -91.6% |    0.08x
    ------------------------------------------------------------

    Total: 325.43ms vs 31.05ms

After:

         N      K |  No OpenMP     OpenMP |    Diff |  Speedup
    ------------------------------------------------------------
       512   2880 |     1.49ms    263.5us |  -82.3% |    0.18x
      2880   4096 |     1.55ms    261.7us |  -83.1% |    0.17x
    201088   2880 |    24.03ms    21.61ms |  -10.1% |    0.90x
    ------------------------------------------------------------

    Total: 78.97ms vs 31.05ms

Tested with unsloth/gpt-oss-20b-GGUF:Q4_K_M.

Signed-off-by: Adrien Gallouët <redacted>
5 weeks agokleidiai : add sme fp16 compute path for q4_0 gemm on aarch64 (llama/20043)
Charles Xu [Tue, 3 Mar 2026 09:40:26 +0000 (10:40 +0100)]
kleidiai : add sme fp16 compute path for q4_0 gemm on aarch64 (llama/20043)

5 weeks agoopencl: add optimized q4_1 mm kernel for adreno (llama/19840)
shaofeiqi [Tue, 3 Mar 2026 03:49:41 +0000 (19:49 -0800)]
opencl: add optimized q4_1 mm kernel for adreno (llama/19840)

* Add Q4_1 OpenCL Kernels

* opencl: refactor transpose

* opencl: format

* opencl: refactor q4_1 unpack

* opencl: move `ggml_cl_mul_mat_q4_1_f32_adreno`

* opencl: refactor `ggml_cl_mul_mat_q4_1_f32_adreno` and kernels

* opencl: rename kernel files and kernes

* opencl: fix build for non adreno

* opencl: move code around and format

---------

Co-authored-by: Li He <redacted>
5 weeks agoggml webgpu: fix workgroup dispatch limit for large batch sizes (llama/19965)
Abhijit Ramesh [Tue, 3 Mar 2026 03:35:11 +0000 (19:35 -0800)]
ggml webgpu: fix workgroup dispatch limit for large batch sizes (llama/19965)

* ggml-webgpu: fix workgroup dispatch limit for large batch sizes

WebGPU limits workgroup sizes to 65535 per dimension. Large MUL_MAT
operations with batch sizes exceedeing this limi would fail.

* add compute_2d_workgroups() helper to split total workgroup ID across
X/Y dimensions

* update mul_mat_reg_tile.wgsl to reconstruct linear workgroup ID from 2D
   dispatch

* update mul_mat_subgroup_matrix.wgsl to reconstruct linear workgroup ID
  from 2D dispatch

* update mul_mat.wgsl to compute global index from 2D workgroup
  coordinates

* refactor all three mul_mat dispatch paths to use the shared helper

* ggml-webgpu: add bounds checking for over-dispatched workgroups

2D workgroup dispatch can over-dispatch when total workgroups don't
divide evenly into the 65535 per-dimension limit. Extra workgroups
would compute invalid batch indices, causing memory corruption.

* add batch_idx bound check to mul_mat_reg_tile.wgsl and
mul_mat_subgroup_matrix.wgsl to prevent over-dispatched workgroups
from accessing invalid memory

* fixes test failures with large batch sizes (eg., bs=[128, 1024])

* ggml-webgpu: add back TODO for spliting large sizes into batches

* Optimize 2d workgroup provisioning

* Set some parameters that increase speed

---------

Co-authored-by: Reese Levine <redacted>
5 weeks agoggml webgpu: Clean up per-thread parameter buffer pool and job submission logic ...
Nikhil Jain [Mon, 2 Mar 2026 18:23:34 +0000 (10:23 -0800)]
ggml webgpu: Clean up per-thread parameter buffer pool and job submission logic (llama/19772)

* Allow webgpu_buf_pool to resize if needed, remove inflight_threads, and replace inflight_threads with num_kernels for submission

* Run clang-format

* Keep track of num batched kernels that have not been submitted yet

* Run clang-format

* Increase buf pool max size

* Increase param buf pool init size

* Remove webgpu buf pool resizing

* Merge with master

* Add buffer pool growth

* Move buffer pool growth outside of lock

* Reduce max pool size to 32

* Run clang-format

* Only resize param buf pool

5 weeks agoggml-webgpu: Support non-contiguous `src0` and overlapping `src0/src1` in binary...
Masashi Yoshimura [Mon, 2 Mar 2026 15:59:53 +0000 (00:59 +0900)]
ggml-webgpu: Support non-contiguous `src0` and overlapping `src0/src1` in binary ops (llama/19850)

* ggml-webgpu: Add binary op support for overlapping and non-contiguous.

* Add newline to binary.wgsl

* Append the test of binary op for src overlapping  to test_bin_bcast.

* Remove unnecessary newline.

5 weeks agovulkan: tune MMVQ for Intel Windows (llama/19988)
Ruben Ortlam [Mon, 2 Mar 2026 14:58:25 +0000 (15:58 +0100)]
vulkan: tune MMVQ for Intel Windows (llama/19988)

5 weeks agoggml-cpu: optimise s390x multiply extend instructions (llama/20032)
Aaron Teo [Mon, 2 Mar 2026 08:23:56 +0000 (16:23 +0800)]
ggml-cpu: optimise s390x multiply extend instructions (llama/20032)

5 weeks agovulkan: improve partial offloading performance on AMD (llama/19976)
Ruben Ortlam [Sun, 1 Mar 2026 16:32:14 +0000 (17:32 +0100)]
vulkan: improve partial offloading performance on AMD (llama/19976)

* vulkan: fix and enable cpy_tensor_async function

* use transfer_queue for async transfers on AMD, synchronize with timeline semaphore

* update offload_op logic

* fix missing transfer submission

* disable async transfer queue on AMD GCN

* revert op batch size change

* fix cpy_tensor_async checks

5 weeks agocuda: cap grid.y at 65535 in non-contiguous dequantize/convert kernels (llama/19999)
oobabooga [Sun, 1 Mar 2026 05:40:22 +0000 (02:40 -0300)]
cuda: cap grid.y at 65535 in non-contiguous dequantize/convert kernels (llama/19999)

5 weeks agoCUDA: add CDNA3 MFMA support for flash attention MMA kernel (llama/19806)
Jayant Lohia [Fri, 27 Feb 2026 18:37:26 +0000 (00:07 +0530)]
CUDA: add CDNA3 MFMA support for flash attention MMA kernel (llama/19806)

* CUDA: add CDNA3 MFMA support for flash attention MMA kernel

Add MI300X (gfx942) MFMA tensor core flash attention using
v_mfma_f32_16x16x16_f16 (FP16 in, FP32 accumulate).

- Add FATTN_WARP_SIZE=64 for CDNA wavefront64
- Add CDNA config for head sizes 64, 80, 96, 112, 128
- Add FP16 MFMA intrinsic path in mma.cuh
- Add manual V transpose load for MFMA register layout
- Route CDNA to MMA for prompt processing, VEC for token generation
- Fix Q loading and combine stride granularity for non-power-of-2 heads

Benchmarks (Qwen2.5-1.5B Q4_K_M, MI300X):
  pp512  +7%,  pp1024 +13%,  pp2048 +23%,  pp4096 +39%
  tg128  -10% (FA overhead, VEC used for both)

All 2480 flash attention tests pass.

Ref: https://github.com/ggml-org/llama.cpp/issues/17917

* address review: replace FATTN_WARP_SIZE with constexpr, improve dispatch

- Replace #define FATTN_WARP_SIZE with constexpr int warp_size =
  ggml_cuda_get_physical_warp_size() in each device function
- Use ne[1]*gqa_ratio threshold for MMA vs tile dispatch. Benchmarked
  crossover on MI300X @ d32768 with power-of-2 GQA models:
    hsk=64  (Llama 1B, gqa=4): MMA wins at eff >= 128 (+11%)
    hsk=128 (Llama 3B, gqa=4): MMA wins at eff >= 128 (+4%)
  Unified threshold: eff_nq >= 128 for all head sizes.
- Remove VEC fallback; small batches fall through to tile kernel

* Update ggml/src/ggml-cuda/fattn.cu

* use ggml_cuda_info().devices warp_size instead of hardcoded check

---------

Co-authored-by: Johannes Gäßler <redacted>
5 weeks agoggml-cpu: add repack for mxfp4 (llama/19738)
Aman Gupta [Fri, 27 Feb 2026 10:15:09 +0000 (18:15 +0800)]
ggml-cpu: add repack for mxfp4 (llama/19738)

7 weeks agoruby : null-check (#3689)
KITAITI Makoto [Thu, 5 Mar 2026 05:36:42 +0000 (14:36 +0900)]
ruby : null-check (#3689)

* Introduce null-check to prevent SEGV

* Fix error message

8 weeks agogguf : sync (ggml/0)
Georgi Gerganov [Fri, 27 Feb 2026 10:24:59 +0000 (12:24 +0200)]
gguf : sync (ggml/0)

8 weeks agoscripts : sync gguf
Georgi Gerganov [Fri, 27 Feb 2026 10:24:33 +0000 (12:24 +0200)]
scripts : sync gguf

8 weeks agotalk-llama : sync llama.cpp
Georgi Gerganov [Fri, 27 Feb 2026 10:23:40 +0000 (12:23 +0200)]
talk-llama : sync llama.cpp

8 weeks agosync : ggml
Georgi Gerganov [Fri, 27 Feb 2026 10:19:27 +0000 (12:19 +0200)]
sync : ggml

8 weeks agoreplace the magic nunber 768 by max work group size to support iGPU (llama/19920)
Neo Zhang [Fri, 27 Feb 2026 01:26:07 +0000 (09:26 +0800)]
replace the magic nunber 768 by max work group size to support iGPU (llama/19920)

Co-authored-by: Neo Zhang Jianyu <redacted>
8 weeks agoggml-zendnn: update code for latest ZenDNN API (llama/19923)
Vishal Singh [Fri, 27 Feb 2026 00:43:41 +0000 (06:13 +0530)]
ggml-zendnn: update code for latest ZenDNN API (llama/19923)

- adapt ggml-zendnn.cpp to the new lowoha::matmul interface
- update the ZenDNN git tag in CMake to the latest release (ZenDNN‑2026‑WW08)
- add static lib support in CMake

8 weeks agoggml : fix AMX and add batched support (llama/19925)
Adrien Gallouët [Thu, 26 Feb 2026 20:39:11 +0000 (21:39 +0100)]
ggml : fix AMX and add batched support (llama/19925)

llama-perplexity -hf ggml-org/Qwen3-0.6B-GGUF:Q4_0 -f wikitext-2-raw/wiki.test.raw -c 2048 -b 2048 --chunks 2

before this commit:

```
perplexity: calculating perplexity over 2 chunks, n_ctx=2048, batch_size=2048, n_seq=1
perplexity: 2.31 seconds per pass - ETA 0.07 minutes
[1]17.3868,[2]22.2199,
Final estimate: PPL = 22.2199 +/- 1.59692

llama_perf_context_print:        load time =     878.56 ms
llama_perf_context_print: prompt eval time =    2037.82 ms /  4096 tokens (    0.50 ms per token,  2009.99 tokens per second)
llama_perf_context_print:        eval time =       0.00 ms /     1 runs   (    0.00 ms per token,      inf tokens per second)
llama_perf_context_print:       total time =    6403.17 ms /  4097 tokens
llama_perf_context_print:    graphs reused =          0
llama_memory_breakdown_print: | memory breakdown [MiB] | total   free    self   model   context   compute    unaccounted |
llama_memory_breakdown_print: |   - Host               |                  845 =   318 +     224 +     302                |
llama_memory_breakdown_print: |   - CPU_REPACK         |                  288 =   288 +       0 +       0                |
llama_memory_breakdown_print: |   - AMX                |                   31 =    31 +       0 +       0                |
```

after this commit:

```
perplexity: calculating perplexity over 2 chunks, n_ctx=2048, batch_size=2048, n_seq=1
perplexity: 1.98 seconds per pass - ETA 0.05 minutes
[1]17.2005,[2]21.8220,
Final estimate: PPL = 21.8220 +/- 1.56485

llama_perf_context_print:        load time =     719.23 ms
llama_perf_context_print: prompt eval time =    1676.23 ms /  4096 tokens (    0.41 ms per token,  2443.58 tokens per second)
llama_perf_context_print:        eval time =       0.00 ms /     1 runs   (    0.00 ms per token,      inf tokens per second)
llama_perf_context_print:       total time =    4258.74 ms /  4097 tokens
llama_perf_context_print:    graphs reused =          0
llama_memory_breakdown_print: | memory breakdown [MiB] | total   free    self   model   context   compute    unaccounted |
llama_memory_breakdown_print: |   - Host               |                  845 =   318 +     224 +     302                |
llama_memory_breakdown_print: |   - AMX                |                  319 =   319 +       0 +       0                |
```
(no more CPU_REPACK)

after this commit, disabling amx:

```
perplexity: calculating perplexity over 2 chunks, n_ctx=2048, batch_size=2048, n_seq=1
perplexity: 2.34 seconds per pass - ETA 0.07 minutes
[1]17.2005,[2]21.8220,
Final estimate: PPL = 21.8220 +/- 1.56485

llama_perf_context_print:        load time =     841.91 ms
llama_perf_context_print: prompt eval time =    2057.28 ms /  4096 tokens (    0.50 ms per token,  1990.98 tokens per second)
llama_perf_context_print:        eval time =       0.00 ms /     1 runs   (    0.00 ms per token,      inf tokens per second)
llama_perf_context_print:       total time =    6454.51 ms /  4097 tokens
llama_perf_context_print:    graphs reused =          0
llama_memory_breakdown_print: | memory breakdown [MiB] | total   free    self   model   context   compute    unaccounted |
llama_memory_breakdown_print: |   - Host               |                  845 =   318 +     224 +     302                |
llama_memory_breakdown_print: |   - CPU_REPACK         |                  319 =   319 +       0 +       0                |
```
=> same perplexity.

Signed-off-by: Adrien Gallouët <redacted>
8 weeks agovulkan: fix fp16 Flash Attention on Windows AMD RDNA2 and below (llama/19921)
Ruben Ortlam [Thu, 26 Feb 2026 18:11:04 +0000 (19:11 +0100)]
vulkan: fix fp16 Flash Attention on Windows AMD RDNA2 and below (llama/19921)

8 weeks agoggml-virtgpu: improve the reliability of the code (llama/19846)
Kevin Pouget [Thu, 26 Feb 2026 12:00:57 +0000 (13:00 +0100)]
ggml-virtgpu: improve the reliability of the code (llama/19846)

* ggml-virtgpu-backend: validate the consistency of the received objects

This patch adds consistency checks in the
ggml-virtgpu-backend (running on the host side) to ensure that the
data received from the guest is consistent (valid pointers, valid
sizes and offsets).

* ggml-virtgpu-backend: add fallback/skips for optional ggml backend methods

```
  1. bck->iface.synchronize(bck)
  2. buft->iface.get_alloc_size(buft, op)
  3. buft->iface.get_max_size(buft)
```

these three methods are optional in the GGML interface. `get_max_size`
was already properly defaulted, but `backend sychronize` and `butf
get_max_size` would have segfaulted the backend if not implemented.

* ggml-virtgpu-backend: fix log format missing argument

* ggml-virtgpu-backend: improve the abort message

* ggml-virtgpu-backend: more safety checks

* ggml-virtgpu-backend: new error code

* ggml-virtgpu-backend: initialize all the error codes

* ggml-virtgpu: add a missing comment generated by the code generator

* ggml-virtgpu: add the '[virtgpu]' prefix to the device/buffer names

* ggml-virtgpu: apir_device_buffer_from_ptr: improve the error message

* ggml-virtgpu: shared: make it match the latest api_remoting.h of Virglrenderer APIR

(still unmerged)

* ggml-virtgpu: update the code generator to have dispatch_command_name in a host/guest shared file

* ggml-virtgpu: REMOTE_CALL: fail if the backend returns an error

* docs/backend/VirtGPU.md: indicate that the RAM+VRAM size is limed to 64 GB with libkrun

* ggml-virtgpu: turn off clang-format header ordering for some of the files

Compilation breaks when ordered alphabetically.

* ggml-virtgpu: clang-format

* ggml-virtgpu/backend/shared/api_remoting: better comments for the APIR return codes

8 weeks agosupport permuted, remove check s0/s10 (llama/19889)
Neo Zhang [Thu, 26 Feb 2026 02:27:20 +0000 (10:27 +0800)]
support permuted, remove check s0/s10 (llama/19889)

Co-authored-by: Neo Zhang Jianyu <redacted>
8 weeks agovulkan: check for memory overlap before doing fusion (llama/19768)
Jeff Bolz [Wed, 25 Feb 2026 17:25:38 +0000 (11:25 -0600)]
vulkan: check for memory overlap before doing fusion (llama/19768)

* vulkan: check for memory overlap before doing fusion

* Update ggml/src/ggml-vulkan/ggml-vulkan.cpp

* address feedback

8 weeks agoggml/gguf : prevent integer overflows (llama/19856)
Georgi Gerganov [Tue, 24 Feb 2026 18:17:11 +0000 (20:17 +0200)]
ggml/gguf : prevent integer overflows (llama/19856)

* gguf : prevent integer overflow for ggml_context mem size

* ggml : fix int overflows in ggml_new_object()

* gguf : prevent string exhaustion

* gguf : prevent array elements exhaustion

* ggml : fix negative tensor type oob

* py : assert that alignment is non-zero power of 2

* ggml : check int overflow in ggml_new_tensor_impl and ggml_new_object

* gguf-py : error on duplicate keys when reading

* py : restore tensor_fields

* enforce proper alignment in add_custom_alignment

* gguf : better name

* gguf : fix ctx size for no_alloc == true

* gguf : minor print fix

* ggml : print values when overflow

* ggml : remove deprecated ggml_type_sizef()

* ggml : relax ggml_type asserts to debug-only

* gguf : add mem_size overflow test

* gguf : add file size check for arrays

* ggml : relax asseerts for ggml_get_type_traits()

* flake8 fix

---------

Co-authored-by: Sigbjørn Skjæret <redacted>
8 weeks agoVulkan Scalar Flash Attention Refactor (llama/19625)
Ruben Ortlam [Tue, 24 Feb 2026 07:35:48 +0000 (08:35 +0100)]
Vulkan Scalar Flash Attention Refactor (llama/19625)

* vulkan: allow using fp16 in scalar flash attention shader

* split rows inside of subgroups for faster synchronization

* use row_split when Br >= 4, change reductions to use shared memory if row_split == 1

* use f32 scalar FA if f16 is not supported by device

* fix amd workgroup size issue

* optimize masksh use

* add medium rows FA shader Br size

* fixes

* add padding to mask shmem buffer

* cache q values into registers for KQ

* fuse lf accumulation, pf and v accumulation into a loop

* stage K loads through shmem

* stage V loads through shmem

* only stage through shmem on Nvidia

* default to Bc 32

* also stage V through shmem when this is done for K

* dynamic subgroups for intel

* use vectorized stores

* use float_type for dequantize4 functions

* use smaller scalar rows size for smaller rows count

* relax flash attention split_k condition to allow non-gqa use

* use minimal subgroup size on Intel

* fix shmem support function

* fix rebase issues

* fixes

* Bc 4 for scalar FA is not a valid configuration

* Use wave32 on AMD RDNA for scalar FA

* add Intel shader core count lookup-table

* fix regressions

* device tuning

* tmpsh size fix

* fix editorconfig

* refactor fa tuning logic into a single place

* fix gqa opt logic

* fix block_rows with small n_rows

* amd tuning

* fix hsk=72/80 issue

* tuning

* allow condition skipping for column check

* use float16 for Of if available

* address feedback

* fix bad RDNA performance on head size <= 128 by limiting occupancy

* allow printing pipeline stats

* cleanup and fixes

* limit occupancy for GCN for small batch FA with large HSK

* disable f16 FA for GCN AMD GPUs on the proprietary driver

8 weeks agovulkan: fix coopmat1 without bf16 support (llama/19793)
Jeff Bolz [Tue, 24 Feb 2026 06:48:32 +0000 (00:48 -0600)]
vulkan: fix coopmat1 without bf16 support (llama/19793)

8 weeks agovulkan: fix data race in mul_mat_id shader (llama/19790)
Jeff Bolz [Tue, 24 Feb 2026 06:43:12 +0000 (00:43 -0600)]
vulkan: fix data race in mul_mat_id shader (llama/19790)

8 weeks agohexagon refactor all Ops to use local context struct (llama/19819)
Max Krasnyansky [Tue, 24 Feb 2026 00:32:14 +0000 (16:32 -0800)]
hexagon refactor all Ops to use local context struct (llama/19819)

* hexagon: refactor set/get/sum-rows ops to use local context

* hexagon: refactor ROPE and Softmax Ops to use local context

Improves performance a bit by precomputing things and saving in the context.

* hexagon: refactor activation ops to use local context struct

* hexagon: refactor unary ops to use local context struct and DMA/VTCM

* hexagon: use aligned hvx_scale function

* hexagon: remove unused fields from op_context

* hexagon: rewrite ROPE to use DMA and VTCM scratchpad

* hex-rope: keep N rows in scratchpad (instead of just two)

* hex-rope: introduce rowidx cache

* hex-rope: remove unused fields

* hex-rope: rewrite dma prefetch logic to allow for multi-row fetch/compute

also removes the need for fastdiv.

* hex-rope: minor formatting

* hex-rope: use indices and unroll the loops

* hex-rope: more updates to cleanup rope-block handling

* hexagon: cleanup supported type/dims checks

* hexagon: all reduce funcs replicated across lanes

There is no need to explicitly replicate the first value.

* snapdragon: update adb and windows scripts to use ubatch-size 256

Updated Ops support handles larger ubatches.

8 weeks agoggml-cpu: arm64: q5_K repack gemm and gemv (and generic) implementations (dotprod...
Alberto Cabrera Pérez [Mon, 23 Feb 2026 12:42:52 +0000 (12:42 +0000)]
ggml-cpu: arm64: q5_K repack gemm and gemv (and generic) implementations (dotprod) (llama/19356)

* Generic GEMV and boilerplate for q5_K dotprod
* Generic GEMM and boilerplate for q5_K dotprod
* ARM64 q5_K dotprod GEMM
* ARM64 q5_K dotprod GEMV

8 weeks agoImprove CUDA graph capture (llama/19754)
Gaurav Garg [Sat, 21 Feb 2026 09:39:36 +0000 (15:09 +0530)]
Improve CUDA graph capture (llama/19754)

* Improve CUDA graph capture

Currently, CUDA graphs are eagerly enabled on the first call to ggml_backend_cuda_graph_compute. If the graph properties keep changing (4+ consecutive updates), the graph is permanently disabled. This is suboptimal because:

- The first call always incurs CUDA graph capture overhead even if the graph is unstable
- Once permanently disabled, CUDA graphs never re-enable even after the graph stabilizes (e.g., switching from prompt processing to decode)

The new approach delays CUDA graph activation until warmup completes: the same cgraph must be called at least twice with matching properties before CUDA graph capture begins. This avoids wasted capture overhead on volatile graphs and allows graphs to become eligible once they stabilize.
This also fixes issues such as https://github.com/ggml-org/llama.cpp/discussions/19708

* Update ggml/src/ggml-cuda/ggml-cuda.cu

Co-authored-by: Johannes Gäßler <redacted>
* Remove EM dashes

* Update ggml/src/ggml-cuda/ggml-cuda.cu

Co-authored-by: Aman Gupta <redacted>
---------

Co-authored-by: Johannes Gäßler <redacted>
Co-authored-by: Aman Gupta <redacted>
8 weeks agoggml-cpu: add RVV vec dot kernels for quantization types (llama/18784)
Taimur Ahmad [Fri, 20 Feb 2026 11:30:07 +0000 (16:30 +0500)]
ggml-cpu: add RVV vec dot kernels for quantization types (llama/18784)

* ggml-cpu: add rvv vec_dot for iq2_s

Co-authored-by: Rehan Qasim <redacted>
* ggml-cpu: add rvv vec_dot for iq3_s

Co-authored-by: Rehan Qasim <redacted>
* ggml-cpu: add rvv vec_dot for tq1_0, tq2_0

Co-authored-by: Rehan Qasim <redacted>
ggml-cpu: add rvv vec_dot for tq1_0, tq2_0

* ggml-cpu: add rvv vec_dot for iq1_s, iq1_m

Co-authored-by: Rehan Qasim <redacted>
* ggml-cpu: add vlen switch for rvv vec_dot

---------

Co-authored-by: Rehan Qasim <redacted>
8 weeks agoggml-webgpu: Add unary op (SQR, SQRT, SIN, COS) support. (llama/19700)
Masashi Yoshimura [Thu, 19 Feb 2026 16:18:30 +0000 (01:18 +0900)]
ggml-webgpu: Add unary op (SQR, SQRT, SIN, COS) support. (llama/19700)

* ggml-webgpu: Add unary op (SQR, SQRT, SIN, COS) support.

* Fix to cast the src value to f32 before sin/cos computing.

8 weeks agovulkan: fix MMQ shader push constants and multi-dispatch (llama/19732)
Ruben Ortlam [Thu, 19 Feb 2026 13:59:16 +0000 (14:59 +0100)]
vulkan: fix MMQ shader push constants and multi-dispatch (llama/19732)

8 weeks agoCUDA: fix kernel selection logic for tile FA (llama/19686)
Johannes Gäßler [Thu, 19 Feb 2026 11:42:58 +0000 (12:42 +0100)]
CUDA: fix kernel selection logic for tile FA (llama/19686)

* CUDA: fix kernel selection logic for tile FA

* add comment

8 weeks agollamafile: powerpc: add FP16 MMA path for Q4/Q8 matmul (llama/19709)
shalinib-ibm [Thu, 19 Feb 2026 06:28:53 +0000 (11:58 +0530)]
llamafile: powerpc: add FP16 MMA path for Q4/Q8 matmul (llama/19709)

Avoid xvi8ger4pp signed→unsigned bias correction by dequantizing Q4/Q8
inputs to FP16 and using FP16×FP16→FP32 MMA. This removes
post-processing overhead and improves performance.

Performance Impact:
1.5 ~ 2x improvement in PP_Speed for Q4 and Q8 Models,
measured with llama-bench and llama-batched-bench.
Q8 Model: granite-4.0-h-micro-Q8_0.gguf (from huggingface)
Q4 Model: Meta-Llama3-8b Q4 model (generated with llama-quantize from
f32 model)

llama-bench Q8 Model Results:
 model                                  size       params   backend      threads              test  Base t/s Patch t/s
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10               pp8           64.48 ± 4.72           73.99 ± 0.27
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10              pp16           80.11 ± 0.32          112.53 ± 0.40
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10              pp32           89.10 ± 0.27          152.95 ± 0.68
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10              pp64           93.65 ± 0.25          187.83 ± 0.83
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10             pp128           99.93 ± 0.02          201.32 ± 0.11
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10             pp256          102.32 ± 0.40          208.32 ± 0.41
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10             pp512          103.42 ± 0.40          209.98 ± 0.14
 granitehybrid 3B Q8_0              3.16 GiB       3.19 B   CPU               10             tg128           20.35 ± 0.01           19.57 ± 0.01

llama-bench Q4 Model Results:
 model                                  size       params   backend      threads              test                Base    t/s                 Patch   t/s
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10               pp8           34.77 ± 0.10           41.23 ± 0.08
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10              pp16           40.81 ± 0.04           64.55 ± 0.15
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10              pp32           44.65 ± 0.05           90.84 ± 0.22
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10              pp64           47.49 ± 0.03          114.39 ± 0.11
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10             pp128           49.29 ± 0.24          120.13 ± 0.19
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10             pp256           49.77 ± 0.23          121.51 ± 0.11
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10             pp512           49.89 ± 0.23          117.52 ± 0.10
 llama 8B Q4_0                      4.33 GiB       8.03 B   CPU               10             tg128           13.40 ± 0.01           13.37 ± 0.00

Llama perplexity Results:

Model                     Base Final PPL Estimate Patch Final PPL Estimate
granite-4.0-h-micro-Q8_0    1.3862 +/- 0.04424         1.3868 +/- 0.04432
Meta-Llama3-8b Q4     1.3801 +/- 0.04116         1.3803 +/- 0.04116

Signed-off-by: Shalini.Salomi.Bodapati <redacted>
8 weeks agoggml webgpu: Fix bug in dispatching large matrix-vector multiplication (llama/19535)
Reese Levine [Wed, 18 Feb 2026 23:06:29 +0000 (16:06 -0700)]
ggml webgpu: Fix bug in dispatching large matrix-vector multiplication (llama/19535)

* Fix bug in dispatching large matrix-vector multiplication

8 weeks agoggml webgpu: shader library organization (llama/19530)
Reese Levine [Wed, 25 Feb 2026 07:33:32 +0000 (09:33 +0200)]
ggml webgpu: shader library organization (llama/19530)

* Basic JIT compilation for mul_mat, get_rows, and scale (ggml/17)

* scale jit working

* preliminary working jit for getrows and mulmat, needs refining

* simplified mul_mat preprocessing switch statement

* get_rows fixes, mul_mat refinement

* formatted + last edits

* removed some extraneous prints

* fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish

* small fix

* some changes, working

* get_rows and mul_mat jit fixed and working

* Update formatting

* formatting

* Add header

---------

Co-authored-by: Neha Abbas <redacted>
Co-authored-by: Reese Levine <redacted>
* Start work on all-encompassing shader library

* refactor argmax, set_rows

* Refactor all but flashattention, mat mul

* flashattention and matrix multiplication moved to new format

* clean up preprocessing

* Formatting

* remove duplicate constants

* Split large shaders into multiple static strings

---------

Co-authored-by: neha-ha <redacted>