sync : ggml#3428

Merged

ggerganov merged 224 commits intomasterfrom

sync-ggml-25-09-20

Sep 29, 2025

Commits on Sep 20, 2025

cuda : remove obsolete sources (ggml/1332)
ggerganov
committed
ggml-cpu: add mxfp4 VSX intrinsics for Power9+ (ppc64le) hardware (llama/15385)

mgiessing
and
ggerganov
committed
musa: handle __hgt2_mask, available starting from MUSA SDK rc4.3.0 (llama/15413)

yeahdongcn
authored and
ggerganov
committed
CANN: optimize rope operator (llama/15335)

YangShuai52
authored and
ggerganov
committed
opencl: mark `argsort` unsupported if cols exceed workgroup limit (llama/15375)

lhez
authored and
ggerganov
committed
musa: fix build warnings (llama/15258)

yeahdongcn
authored and
ggerganov
committed
vulkan: shorten pipeline name strings (llama/15431)

jeffbolznv
authored and
ggerganov
committed
CUDA: replace GGML_CUDA_F16 with CUDA arch checks (llama/15433)

JohannesGaessler
authored and
ggerganov
committed
CUDA: refactor FA support/selection code (llama/15454)

JohannesGaessler
authored and
ggerganov
committed
sched : copy only the used experts when offloading prompt processing (llama/15346)

slaren
authored and
ggerganov
committed
musa: add GGML_UNUSED_VARS (llama/15446)

yeahdongcn
authored and
ggerganov
committed
ggml : fix condition of im2col on Metal backend (llama/15460)

ngxson
authored and
ggerganov
committed
vulkan: Reuse conversion results in prealloc_y (llama/15410)

jeffbolznv
authored and
ggerganov
committed
vulkan: add exp operation (llama/15456)

authored and
ggerganov
committed
vulkan : support conv_2d_dw with f16 weights (llama/15392)

Acly
authored and
ggerganov
committed
sched : fix possible use of wrong ids tensor when offloading moe prompt processing (llama/15488)

slaren
authored and
ggerganov
committed
CANN: Optimize RMS_NORM using cache (llama/15419)

noemotiovon
authored and
ggerganov
committed
ggml-cpu: Support Q5_0 and Q5_1 on s390x (llama/15486)

taronaeo
authored and
ggerganov
committed
cuda : add Pad Reflect 1D support (llama/14659)

authored and
ggerganov
committed
ggml: add `conv3d` op (llama/15182)

rmatif
authored and
ggerganov
committed
ggml WebGPU: add support for quantization types (llama/15440)

reeselevine
authored and
ggerganov
committed
vulkan: optimize mul_mat_id loading row ids into shared memory (llama/15427)

jeffbolznv
authored and
ggerganov
committed
vulkan : support ggml_mean (llama/15393)

Acly
authored and
ggerganov
committed
vulkan: Rewrite synchronization to allow some overlap between nodes (llama/15489)

jeffbolznv
authored and
ggerganov
committed
vulkan: optimize rms_norm, and allow the work to spread across multiple SMs (llama/15281)

jeffbolznv
authored and
ggerganov
committed
CUDA: fix half2 -> half conversion for HIP (llama/15529)

JohannesGaessler
authored and
ggerganov
committed
vulkan: workaround MoltenVK compile failure in multi_add (llama/15506)

authored and
ggerganov
committed
vulkan: enable Conv2D for Apple after MoltenVK fixed the bug (llama/15526)

0cc4m
authored and
ggerganov
committed
vulkan: Support FA with any multiple of 8 head sizes (llama/15537)

jeffbolznv
authored and
ggerganov
committed
vulkan: apply MUL_MAT_ID subgroup optimization to non-coopmat devices (llama/15524)

0cc4m
authored and
ggerganov
committed
CANN: ROPE cache sin/cos repeat (llama/15501)

noemotiovon
authored and
ggerganov
committed
metal : add FA kernels for HS=40 (llama/15559)
ggerganov
committed
CUDA: MoE helper in device code, better tile sizes (llama/15525)

JohannesGaessler
authored and
ggerganov
committed
metal: fix regression when no metal devices are present (llama/15531)

booxter
authored and
ggerganov
committed
vulkan: fix min subgroup 16 condition for mmid subgroup optimization (llama/15565)

0cc4m
authored and
ggerganov
committed
opencl: fix support ops condition for `rms_norm` (llama/15560)

lhez
authored and
ggerganov
committed
CUDA: Accelerate MXFP4 table lookup using `__byte_perm` (llama/15451)

authored and
ggerganov
committed
vulkan: Remove splitting for mul_mat_id (llama/15568)

jeffbolznv
authored and
ggerganov
committed
Add a warning for special devices (llama/15563)

authored and
ggerganov
committed
metal : remove contiguous assertion for src0 in IM2COL (llama/15577)

CISC
authored and
ggerganov
committed
metal : improve `MUL_MAT_ID` (llama/15541)
ggerganov
committed
metal : optimize FA vec for large sequences and BS <= 8 (llama/15566)
ggerganov
committed
CUDA: return -1 for nonexistent compiled arch (llama/15587)

JohannesGaessler
authored and
ggerganov
committed
llamafile: PowerPC Sgemm Optimization (llama/15558)

shalinib-ibm
authored and
ggerganov
committed
SYCL: fix rms_norm_mul_add for tensor dim not a multiple of sg_size (llama/15592)

qnixsynapse
authored and
ggerganov
committed
OpenCL: add fused group_norm/norm, mul, add (llama/15314)

rmatif
authored and
ggerganov
committed
ggml-cpu : add basic RVV support for vector f32 ops (llama/15057)

xctan
authored and
ggerganov
committed
CANN: refactor mask handling and improve performance in FA (llama/15561)

noemotiovon
authored and
ggerganov
committed
HIP: Enable support for ggml_backend_cuda_register_host_buffer (llama/15615)

IMbackK
authored and
ggerganov
committed
cuda: Add cublasLt_static linking when GGML_STATIC is enabled (llama/15622)

matiaslin
authored and
ggerganov
committed
kv-cache : remove LLAMA_SET_ROWS checks (llama/15505)
ggerganov
committed
ggml : fix SSM_SCAN for n_groups > 1 (llama/15625)

compilade
authored and
ggerganov
committed
ggml-cpu: fix invalid hsum build in debug s390x (llama/15634)

taronaeo
authored and
ggerganov
committed
CUDA: add conv2d (llama/15635)

mnehete32
authored and
ggerganov
committed
CUDA: fuse adds, fuse add with rms norm (llama/15631)

am17an
authored and
ggerganov
committed
CUDA: fix bug in rms_norm fusion (llama/15660)

am17an
authored and
ggerganov
committed
CANN: FIx compiler warnings (llama/15661)

noemotiovon
authored and
ggerganov
committed
vulkan: Skip syncing for prealloc_y when it is reused (llama/15544)

jeffbolznv
authored and
ggerganov
committed
CUDA: use FP32 arithmetic for conv2d (llama/15683)

JohannesGaessler
authored and
ggerganov
committed
llama: use FA + max. GPU layers by default (llama/15434)

JohannesGaessler
authored and
ggerganov
committed
ggml: update kleidiai to v1.13.0 (llama/15663)

chaxu01
authored and
ggerganov
committed
vulkan: clamp matmul and FA results to the max finite value (llama/15652)

jeffbolznv
authored and
ggerganov
committed
vulkan: Allow fallback to sysmem memory when vidmem is full (llama/15649)

jeffbolznv
authored and
ggerganov
committed
vulkan : remove unused portability_enumeration_ext variable (llama/15679)

danbev
authored and
ggerganov
committed
vulkan: mul_mat_id coopmat2 optimizations (llama/15546)

jeffbolznv
authored and
ggerganov
committed
vulkan: handle large sizes for get_rows (llama/15686)

jeffbolznv
authored and
ggerganov
committed
llama : separate compute buffer reserve from fattn check (llama/15696)

slaren
authored and
ggerganov
committed
metal : fix checks for available FA kernels (llama/15700)
ggerganov
committed
CANN: fix RoPE cache issue on multi-device (llama/15629)

hipudding
authored and
ggerganov
committed
CANN: Optimize MUL_MAT_ID (llama/15658)

hipudding
authored and
ggerganov
committed
CUDA: fix build error from ambiguous __half conversions in conv2d (llama/15690)

qnixsynapse
authored and
ggerganov
committed
ggml : WebGPU add TRANSPOSE and RESHAPE to supported ops (llama/15695)

authored and
ggerganov
committed
Vulkan: Add Integer Dot Product mul_mat_vec shader for legacy quants (llama/14903)

0cc4m
authored and
ggerganov
committed
ggml: aarch64: Implement SVE F16 kernels for vector functions (llama/15115)

authored and
ggerganov
committed
ggml: SVE support for exponential functions (llama/15145)

s-goto-11
and
ggerganov
committed
vulkan: disable large mmv subgroups on older Nvidia GPUs (llama/15717)

0cc4m
authored and
ggerganov
committed
vulkan: add missing clamps in new mul_mat_id paths (llama/15702)

jeffbolznv
authored and
ggerganov
committed
vulkan: use memory budget extension to read memory usage (llama/15545)

giladgd
authored and
ggerganov
committed
ggml-backend: raise GGML_MAX_SPLIT_INPUTS (llama/15722)

JohannesGaessler
authored and
ggerganov
committed
CANN: Support ext_factor in rope (llama/15710)

hipudding
authored and
ggerganov
committed
CANN: Support eager execution mode under ACL graph compilation (llama/15712)

noemotiovon
authored and
ggerganov
committed
opencl: add attn sinks support for FA kernels (llama/15706)

rmatif
authored and
ggerganov
committed
vulkan: Fix macro parameter order for f32 matmul shaders (llama/15716)

jeffbolznv
authored and
ggerganov
committed
CANN: Resolve soft_max precision issue (llama/15730)

hipudding
authored and
ggerganov
committed
vulkan: fix shaders gen when no integer dot is available (llama/15740)

0cc4m
authored and
ggerganov
committed
CANN: Fix type float_t to float (llama/15736)

noemotiovon
authored and
ggerganov
committed
CANN: Mask unsupported TRANSPOSE_1D operator (llama/15733)

hipudding
authored and
ggerganov
committed
ggml-cpu : optimize RVV kernels (llama/15720)

xctan
authored and
ggerganov
committed
CANN: Add RoPE contiguous check for 310I DUP device (llama/15735)

hipudding
authored and
ggerganov
committed
CUDA: Optimize `rms_norm_f32` kernel and its fused variants, giving 1-6% perf E2E (llama/15715)

authored and
ggerganov
committed
ggml vulkan: add hardsigmoid and hardswish operations (llama/15762)

relent95
authored and
ggerganov
committed
vulkan : update ggml_vk_instance_validation_ext_available (llama/15666)

danbev
authored and
ggerganov
committed
vulkan: don't use std::string in load_shaders, to improve compile time (llama/15724)

jeffbolznv
authored and
ggerganov
committed
vulkan: fix mmv subgroup16 selection (llama/15775)

0cc4m
authored and
ggerganov
committed
CANN: fix acl_rstd allocation size in ggml_cann_rms_norm (llama/15760)

authored and
ggerganov
committed
opencl: add hs=40 to FA (llama/15758)

rmatif
authored and
ggerganov
committed
CANN: Fix precision issue on 310I DUO multi-devices (llama/15784)

hipudding
authored and
ggerganov
committed
ggml: add ops for WAN video model (cuda && cpu) (llama/15669)

authored and
ggerganov
committed
CANN: Refactor ND to NZ workspace to be per-device (llama/15763)

authored and
ggerganov
committed
metal : Add template specialization for mul_mm_id w/ ne20 == 10 (llama/15799)

gabe-l-hart
authored and
ggerganov
committed
CUDA : conditionally add cuda architectures (ggml/1341)

gjasny
authored and
ggerganov
committed
ggml : introduce semantic versioning (ggml/1336)

danbev
authored and
ggerganov
committed
CUDA: fastdiv, launch bounds for mmvq + q8_1 quant (llama/15802)

JohannesGaessler
authored and
ggerganov
committed
ggml-cpu: drop support for nnpa intrinsics (llama/15821)

taronaeo
authored and
ggerganov
committed
ggml-cpu: document use of "free" memory [no ci] (llama/15834)

JohannesGaessler
authored and
ggerganov
committed
kleidiai: generalize compute_forward_kv_cache to compute_forward_fp16 (llama/15817)

chaxu01
authored and
ggerganov
committed
CUDA: faster tile FA (Pascal/AMD), headsize 256 (llama/15769)

JohannesGaessler
authored and
ggerganov
committed
ggml WebGPU: remove userdata from request adapter callback (llama/15527)

danbev
authored and
ggerganov
committed
vulkan: Use larger loads in scalar/coopmat1 matmul (llama/15729)

jeffbolznv
authored and
ggerganov
committed
vulkan: Support pad_ext (llama/15794)

jeffbolznv
authored and
ggerganov
committed
ggml-cpu: clean up s390x SIMD (llama/15855)

taronaeo
authored and
ggerganov
committed
vulkan: support im2col_3d (llama/15795)

jeffbolznv
authored and
ggerganov
committed
CANN: Stream sync between devices for acl_graph (llama/15809)

authored and
ggerganov
committed
CUDA: non-contiguous src0 not supported for PAD (llama/15869)

CISC
authored and
ggerganov
committed
ggml: allow casting between f32 and i32 (llama/15783)

ngxson
authored and
ggerganov
committed
metal : refactor + optimize (llama/15857)
ggerganov
committed
cuda : fix supports_op condition for get_rows when number of blocks is too large (llama/15868)
ggerganov
committed
CUDA: generate_cu_files.py - add missing mxfp4 (llama/15880)

am17an
authored and
ggerganov
committed
vulkan: sort graph to allow more parallel execution (llama/15850)

jeffbolznv
authored and
ggerganov
committed
CUDA: fix GET_ROWS for large tensors (llama/15882)

JohannesGaessler
authored and
ggerganov
committed
CUDA: Add mul_mat_id support for the mmf kernel (llama/15767)

am17an
authored and
ggerganov
committed
Workaround for subgroup arithmetic failing on MoltenVK with AMD GPUs (issue 15846) (llama/15886)

lksj92hs
authored and
ggerganov
committed
HIP: use v_dot2_f32_f16 instruction for FA (llama/15884)

JohannesGaessler
authored and
ggerganov
committed
vulkan: Fix OOB accesses in soft_max_back (llama/15861)

jeffbolznv
authored and
ggerganov
committed
vulkan: throw the oom error instead of no memory type found (llama/15905)

0cc4m
authored and
ggerganov
committed
CANN: implement LRU cache for ACL graphs (llama/15814)

noemotiovon
authored and
ggerganov
committed
CANN: Add ROPE sin/cos cache for reuse (llama/15912)

authored and
ggerganov
committed
sync : ggml
ggerganov
committed
metal : make the backend async (llama/15906)
ggerganov
committed
sync : ggml
ggerganov
committed
ggml-cpu : fix padding in ggml_timestep_embedding (llama/15917)

danbev
authored and
ggerganov
committed
CUDA: Add `fastdiv` to `k_bin_bcast*`, giving 1-3% E2E performance (llama/15872)

ORippler
authored and
ggerganov
committed
CANN: Disable acl_graph for prefill stage (llama/15933)

hipudding
authored and
ggerganov
committed
kleidiai: fix GGML_ASSERT(*cur_backend_id != -1) failed (llama/15614)

chaxu01
authored and
ggerganov
committed
ggml-cpu : add check for ARM MATMUL_INT8/i8mm support (llama/15922)

danbev
authored and
ggerganov
committed
CUDA: larger SRAM reads for tile FA, AMD FP16 dot (llama/15927)

JohannesGaessler
authored and
ggerganov
committed
ggml-backend : add GGML_BACKEND_DEVICE_TYPE_IGPU device type (llama/15797)

slaren
authored and
ggerganov
committed
Revert "sycl: add usage of enqueue_functions extension (llama/14244)" (llama/15910)

NeoZhangJianyu
authored and
ggerganov
committed
vulkan: Make device memory check more portable (llama/15939)

mbaudier
authored and
ggerganov
committed
Vulkan iGPU device selection overhaul and PCI ID API support (llama/15947)

0cc4m
authored and
ggerganov
committed
ggml-zdnn: fix #15414, activate FP16 and BF16 acceleration and incorrect zTensor free (llama/15839)

taronaeo
authored and
ggerganov
committed
metal : fix memory leaks (llama/15962)
ggerganov
committed
metal : allow ops to run concurrently (llama/15929)
ggerganov
committed
metal : refactor kernel loading (llama/15964)
ggerganov
committed
vulkan: initialize vulkan-hpp to allow using extension function pointers (llama/15705)

jeffbolznv
authored and
ggerganov
committed
vulkan: fix failing dequant shaders (llama/15862)

jeffbolznv
authored and
ggerganov
committed
ggml-zdnn: rm user mapped buffers (llama/15965)

taronaeo
authored and
ggerganov
committed
metal : fix kernel requirements (llama/15983)
ggerganov
committed
Vulkan: Clean up mul_mm shader (llama/15987)

0cc4m
authored and
ggerganov
committed
metal : remove memory pools (llama/15966)
ggerganov
committed
CUDA: some micro-optimizations in mmf.cuh for mul_mat_id (llama/15926)

am17an
authored and
ggerganov
committed
SYCL: Add COUNT_EQUAL operator support (llama/15991)

yael-works
authored and
ggerganov
committed
CUDA: fix im2col_3d to respect non-contiguous inputs (views) (llama/15956)

authored and
ggerganov
committed
ggml : fix padding in timestep embedding kernels (llama/15932)

danbev
authored and
ggerganov
committed
CANN: Optimize ggml_cann_set_device (llama/15935)

noemotiovon
authored and
ggerganov
committed
vulkan: automatically remove unsupported devices (llama/15976)

netrunnereve
authored and
ggerganov
committed
CUDA: fix FA occupancy, optimize tile kernel (llama/15982)

JohannesGaessler
authored and
ggerganov
committed
sync : ggml
ggerganov
committed
metal : refactor + optimize v2 (llama/15995)
ggerganov
committed
GGML WebGPU: Support for ADD, MUL, RMS_NORM, GET_ROWS operators (llama/16018)

authored and
ggerganov
committed
CANN: Remove print (llama/16044)

noemotiovon
authored and
ggerganov
committed
metal : handle nil cv during pipeline creation (llama/16065)
ggerganov
committed
metal : avoid call free for non-owned buffer (llama/16067)

jhen0409
authored and
ggerganov
committed
metal : improve F32, F16 and BF16 mat-vec multiplication (llama/16057)
ggerganov
committed
cuda : add missing F32<->I32 entries in ggml_cuda_cpy_fn (llama/16060)

CISC
authored and
ggerganov
committed
metal : use function constants for mul_mv_ext kernels (llama/16074)
ggerganov
committed
CUDA: fix compilation on CC 6.0 (llama/16091)

JohannesGaessler
authored and
ggerganov
committed
CUDA: Optimize PAD_REFLECT_1D (llama/15957)

authored and
ggerganov
committed
rename optimize_graph to graph_optimize (llama/16082)

jeffbolznv
authored and
ggerganov
committed
opencl: optimize mxfp4 kernels (llama/16037)

authored and
ggerganov
committed
cmake : fix static linking for OpenMP on Unix-like systems (llama/16031)

angt
authored and
ggerganov
committed
ggml-amx : fix ggml_amx_init() on generic Linux (llama/16049)

angt
authored and
ggerganov
committed
ggml : refactor forward_dup for cpu backend (llama/16062)

ngxson
authored and
ggerganov
committed
vulkan: use vec dot for matrix matrix multiplications (llama/16056)

0cc4m
authored and
ggerganov
committed
sync : ggml
ggerganov
committed
talk-llama : sync llama.cpp
ggerganov
committed

Commits on Sep 29, 2025

ggml : bump version to 0.9.1
ggerganov
committed
ggml : prepare for development of 0.9.2-dev
ggerganov
committed
vulkan: fix validation error about VK_PIPELINE_CREATE_CAPTURE_STATISTICS_BIT_KHR (llama/16086)

jeffbolznv
authored and
ggerganov
committed
vulkan: optimize UMA buffer operations and fix driver hangs (llama/16059)

giuseppe
authored and
ggerganov
committed
opencl: initial `q8_0` mv support (llama/15732)

lhez
authored and
ggerganov
committed
opencl: fix concat crash on win arm64 with Adreno (llama/15944)

lhez
authored and
ggerganov
committed
vulkan: vec dot matrix multiplication fix (llama/16151)

0cc4m
authored and
ggerganov
committed
vulkan: add RTE variants of exp shader (llama/16165)

jeffbolznv
authored and
ggerganov
committed
Vulkan: add conv_transpose_2d operation (llama/16022)

relent95
authored and
ggerganov
committed
ggml : add ggml_op_is_empty (llama/16122)
ggerganov
committed
ggml : extend ggml_can_fuse to work with non-sequential nodes (llama/16123)
ggerganov
committed
ggml : implement set_rows with i32 index (llama/16159)

committed
ggml-cpu : fix typo in gemm comments [no ci] (llama/16189)

danbev
authored and
ggerganov
committed
zdnn: refactor codebase + add docs (llama/16178)

taronaeo
authored and
ggerganov
committed
ggml : fix uninitialized is_on_grid in quantize_row_iq3_xxs_impl (llama/15928)

CISC
authored and
ggerganov
committed
ggml-cpu: Respect cpumask settings (llama/16164)

wishstudio
authored and
ggerganov
committed
ggml : split graph allocations according to backend max buffer size (llama/15815)

Acly
authored and
ggerganov
committed
llama: print memory breakdown on exit (llama/15860)

JohannesGaessler
authored and
ggerganov
committed
rpc : use ggml logging facilities

rgerganov
authored and
ggerganov
committed
metal : restore im2col perf (llama/16219)
ggerganov
committed
metal : relax reorder conditions (llama/16216)
ggerganov
committed
metal : fuse NORM + MUL + ADD, support non-multiples of 4 (llama/16220)
ggerganov
committed
ggml : bump version to 0.9.3 (ggml/1353)

danbev
authored and
ggerganov
committed
ggml : remove -dev suffix from release version (ggml/1355)

danbev
authored and
ggerganov
committed
ggml : fix loongarch lsx compilation error (llama/15864)

junchao-loongson
authored and
ggerganov
committed
CUDA: add a fused top-K MoE kernel (llama/16130)

am17an
authored and
ggerganov
committed
musa: fix build warnings (llama/15611)

yeahdongcn
authored and
ggerganov
committed
ggml-cpu: implement MXFP4 SIMD for s390x (llama/16193)

taronaeo
authored and
ggerganov
committed
common : use cpp-httplib as a cURL alternative for downloads (llama/16185)

angt
authored and
ggerganov
committed
metal : report OOM errors (llama/16274)
ggerganov
committed
devops: add s390x & ppc64le CI (llama/15925)

authored and
ggerganov
committed
vulkan: support GET_ROWS for k-quants (llama/16235)

jeffbolznv
authored and
ggerganov
committed
vulkan: throw system error instead of SIGABRT during init on older devices (llama/16156)

DmyMi
authored and
ggerganov
committed
CUDA: refactor and deduplicate vector FA kernels (llama/16208)

JohannesGaessler
authored and
ggerganov
committed
CUDA: mul_mat_id for mmf for bs <= 64 for f16 and bs <= 32 for f32 (llama/16277)

am17an
authored and
ggerganov
committed
vulkan : make the vulkan.hpp dynamic dispatcher instance private (llama/16224)

Acly
authored and
ggerganov
committed
vulkan: support arbitrary KV dimension in flash attention (llama/16160)

jeffbolznv
authored and
ggerganov
committed
vulkan: handle mat_mul with A matrix > 4GB (llama/16176)

jeffbolznv
authored and
ggerganov
committed
metal : fuse non-sequential nodes (llama/16102)
ggerganov
committed
metal : extend mat-mat multiplication support (llama/16225)
ggerganov
committed
vulkan: 64-bit im2col (llama/16135)

jeffbolznv
authored and
ggerganov
committed
ggml : fix GGML_F32_VEC_FMA argument order in ggml_vec_mad1_f32 (llama/16307)

CISC
authored and
ggerganov
committed
vulkan: Fix validation failure in quantized flash attention (llama/16292)

jeffbolznv
authored and
ggerganov
committed
ggml : fix dependencies for ggml_set_rows (llama/16318)
ggerganov
committed
ggml : check cuda and metal argsort limits and add test (llama/16323)

CISC
authored and
ggerganov
committed
cmake : remove metal flag (llama/0)
ggerganov
committed
sync : ggml
ggerganov
committed
talk-llama : sync llama.cpp
ggerganov
committed