125 Commits

Include renames

Author	SHA1	Message	Date
+6878 people	2ca53bb45e	sync : ggml (#3962) ... * hexagon: tiling, tracing and optimizations for unary ops (llama/25474) * hexagon: tile wide rows in pointwise unary ops to avoid VTCM overflow * unary: reject permuted tensors for now (not used by models) * hex-unary: replace divs with fastdiv * hex-unary: add vtcm layout and host computed kernel params * hex-unary: move fastdiv init into kernel params * hex-unary: add specialized thread functions to improve generated code * hex-unary: tracing instrumentation for unary ops * hex-unary: factor out hvx kernels, streamline and remove more duplication * ggml-hexagon: fix std::min collision with Windows min macro * hex-cmake: make lto build happy --------- Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com> * ggml : process data in smaller chunks in CUDA ggml_top_k() and ggml_argsort() to reduce temporary buffers memory usage (llama/24776) * ggml : process data in smaller chunks in CUDA ggml_top_k() implementation to reduce temporary buffers memory usage * ggml : allocate tmp_dst only only once before the loop * chore : whitespaces Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * ggml : use chunked processing in both CUDA CUB top-k and argsort implementations * chore : separate argsort_f32_i32_cuda_bitonic() call from return statement Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * chore : replace ternary operators with min/max --------- Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * opencl: cluster-parallel decode FA for Adreno (llama/25473) * ggml-et: Initial ET backend (llama/24179) * ggml-et: Add performance logging * ggml-et: Quants helpers * ggml-et: Add MUL_MAT kernel * ggml-et: Add ROPE kernel * ggml-et: Add RMS_NORM kernel * ggml-et: Add GLU kernel * ggml-et: Add SOFT_MAX kernel * ggml-et: Add GET_ROWS kernel * ggml-et: Add CONT kernel * ggml-et: Add SET_ROWS kernel * ggml-et: Add MUL_MAT_ID kernel * ggml-et: Build et kernels as part of ggml * ggml-et: Embed kernels with fs fallback * ggml-et: Build fixes * ggml-et: Add MUL_MAT F32xF32 op * ggml_et: Add MUL_MAT_ID op * ggml-et: Disable offloading for debug * ggml-et: Refactor out block ops * ggml-et: ggml backend API changes * ggml-et: Add RESHAPE/TRANSPOSE to supported * ggml-et: Add CONT_F16 * ggml-et: Add supported ops doc * gglm-et: Initial doc * ggml-et: Remove runtime import hacks We can now import the runtime by a simple find_package(), so we can cleanup the CMakeLists.txt. * ggml-et: Fix GET_ROWS kernel Fix lost batch dimension. Also clean vibe-comments. * ggml-et: Fix SET_ROWS kernel Remove incorrect broadcasting guard. * ggml-et: Use custom instruction for fp32->fp16 * ggml-et: Vectorize set_rows fp32->fp16 * ggml-et: Fix ROPE kernel (yarn) ggml-et: fix et_logf WIP: Fix ramp WIP: fix ROPE! * ggml-et: Better sinf * ggml-et: Fix SOFT_MAX Add `max_bias` and `sink` support. * ggml-et: Fix CONT Reorder from contiguous write to read with atomic stores. * ggml-et: Fix elmap kernel Remainder handlin * ggml-et: Fix MUL_MAT MUL_MAT_ID remainders * ggml-et: Fix ET-SOC reference * ggml-et: Fix embed kernels scripts for old python This allows GGML-ET to build on pre-3.8 python. * Add sysemu support with compile time flag `-DGGML_ET_SYSEMU=ON` (llama/6) * Example using ET-Soc-1 emulator configuration Example usage: ```bash cmake -B build -DGGML_CUDA=OFF -DGGML_ET=ON -DLLAMA_CURL=OFF -DGGML_CCACHE=ON cmake --build build --config Release -j $(nproc) time ./build/bin/test-backend-ops ./build/bin/llama-server \ --model Qwen3-0.6B-Q8_0.gguf \ --alias Qwen3-0.6B-Q8_0 \ -fa 0 \ --ctx-size 1024 \ --no-warmup \ --host 127.0.0.1 \ --port 8080 ``` * build: proper dep tracking for kernels * support host using MOLD linker * initial multi core GET_ROW F32 implementation * vectorized q8 dequant * wip: cland warning clenaups and initial logging refactor * wip: message default message cleanup * chore: message cleanups * cmake cleanup * migrate to use platform provided functions * cmake back into subdir * support et_print() in kernels * fix: repair kernel building * perf: operations run async by default * debug: proper kernel dep tracking and error detection on kenrel launch * fix: kernel binary dep tracking and fixing get_rows_f32 erroring * perf: back to doing async kernel runs by default * perf: vectorize and parallel device memset * merge matmul work * misc: align allocation and enable all offload * misc: delete deadcode and respect memory limits * fix: repair tensor debug print * fix: loosen RMS_NORM op percision * feat: Q4_0 GET_ROWS * perf: FP32 MUL_MAT using TensorFMA * update limitations * perf: redue L1 load in compute_block_dot_product_q8_0 * feat: save kernel mapping (name to id) when profiling is enabled * chore: memops cleanup * perf: parallelize softmax by rows * perf: vectorize 2nd phase of softmax * perf: ban GET_ROWS from offloaded * perf: vectorize and non-atomic for eltwise ops and sub support * perf: vectorize normal rope * perf: glu runs in parallel * merge: manually merge saqib's work on kernel fixes * perf: more vectorized RoPE * perf: parallelize mul_mat_id * perf: parallelize set_rows_f32 * perf: vectorize softmax * feat: support kernel fusion and fuse RMS_NORM + MUL * fix: mostly resolve test-backend-ops failure in SOFT_MAX and ROPE * fix: bump max rope dims for gemma * feat: GeGLU and SCALE support to fully offload Gemma * perf: faster device memset * feat: get_rows supporting Q4_K and avoid cont cache coherent issues * better F32 MM * feat: NORM for ET backend * feat: SQR for ET backend * feat: UNARY on ET * feat: el_map support broadcasting for ET * feat: SUM_ROWS in ET backend * feat: more ops in ET backend * feat: WKV* operators in ET backend * perf: parallelize operators across cacheline instead of row * perf: parallelize get_rows on cacheline * wip: baseline FlashAttention for ET backend * wip: enough FA and CPY f32->f16 to run llama 3.1 fully offloaded with FA on * feat: f16 x f16 -> f32 MM using matrix engine * wip: f16 FlashAttention using matrix engine * wip: clean up * feat: barriers * perf: optimize FA_F16 in ET * perf: vectorize pack_k_for_transpose16 * perf: prefetch next loop matrix tile * perf: FlashAttention 2nd MM uses TensorFMA and optimizations * cleanup: flashattention reorg * perf: optimizations and fixes * feat: L2SCP API and make FlashAttention support DV = 256 for gemma * perf: parallelize norms beyond single row * feat: GATED_DELTA_NET support and relaxed L2_NORM requirment * feat: loosen RMS_NORM, NORM, ROPE contingous req too * feat: repeat supports brocasting on dim 0 and loosen cont check * feat: FILL and DIAG operator * feat: loosen UNARY support chcek * feat: TRI support * feat: SOLVE_TRI support * feat: basic SET support * feat: loosen CONT req * perf: fp16_to_fp32 use ASM * feat: IMROPE support * feat: PAD support * feat: global barrier * fix: view must live on the same backend as backing tensor * feat: relax CONCAT in ET backend * feat: dead simple CUMSUM implementation * feat: basic SSM_CONV support * feat: loosen CONCAT req * feat: relax GATED_DELTA_NET and add SET support proper * cleanup: cleanup LCM math * feat: SWIGLU single input * feat: SSM_SCAN support * feat: el_map supports non aligned tensors in best effort * feat: basic GROUP_NORM support * feat: loosen MUL_MAT capablities slightly * feat: loosen MUL_MAT and GET_ROWS and add IM2COL * feat: special case for softmax 1x1x1x1 * feat: loosen SOFT_MAX req in ET backend * fix: el_map unaligned acse fixes * perf: optimize zero_acc_vec in flash_attn_ext_f16_me * perf: use hart 1 for packing in MM and FA for FP16 * feat: kernel semaphore * perf: better instruction sequence in FlashAttention * fix: gated_delta_net with proper masking * perf: better parallelization for GATED_DELTA_NET * perf: parallelize SSM_CONV over nr * perf: vectorize SSM_CONV * perf: optimize MUL_MAT for q8 * feat: support Gemma 4 * fix: support multi-device * feat: broader GLU support * feat: unary ops supports view * fix: repair fp16 MM using matrix engine * perf: handle large N GEMV better * perf: better q8_0 MM * perf: better set_rows * add back deleted files * fix: repair after merge * feat: POC version of uberkernel * feat: RMS_NORM in uberkernel * feat: add more kernels into usage * chore: clean up uberkernel compilation * perf: faster flash attention * perf: opt flash attention for large seq length * feat: loosen op bounds. clamp and mean support * perf: vectorize ssm_scan * perf: slightly faster FA * perf: FlashAttention parallel MM and load * perf: fuse Q8 MM and ADD * feat: basic conv kernel for ET * softMAx_test * set_rows_f32 * get_rows and cont * testing * set_rows_exp * Junk addition * Narrowing the issue * Update flash_attn_ext_f16_me.c Focusing FA_ext_f16_me * test * Eviction updated * Detailed cache eviction debug * mulmat * removeal of `BUILD_FOR_UBERKERNEL` flag * cleaning... * fix: balance FCC0 count * feat: implement mul_mat and mul_mat_id for Q4_0 type * optimize uberkernel plan upload * add mul_mat q4 into uberkernel * enable gating flush to just uberkernel * update docs for ET * update op support for ET * et-backend: optimize Q4_0 and Q8_0 mul_mat_id row accumulations * et-backend: specialize mul_mat_id kernels for Q4_0 and Q8_0 * et-backend: fix RoPE YaRN corr_dim formula and handle degenerate inputs * test-backend-ops: add DeepSeek-V2-Lite RoPE test coverage * et-backend: add Q4_0 mul_mat matrix-engine kernel using TensorFMA32 * et-backend: vectorize Q4_0 matrix-engine dequantization * et-backend: support hybrid matrix/vector engine execution for Q4_0 mul_mat tail * et-backend: run partial-N tiles on matrix engine for Q4_0 mul_mat * et-backend: route Q4_0 mul_mat N < 53 to vecdot for better prefill latency * Update uberkernel.c * Update unary_f32.c * gemma 4 * bisect gemma4: enable scale_f32 only * bisect gemma4: +rms_norm_f32 * bisect gemma4: +rms_norm_mul_f32 * bisect gemma4: disable rms_norm_mul_f32 -- BREAKS OUTPUT * bisect gemma4: +rope_f32 (skip rms_norm_mul) * bisect gemma4: +el_map_f32 * bisect gemma4: +softmax_f32 * bisect gemma4: +get_rows_f32 * bisect gemma4: +glu_f32 * bisect gemma4: +mul_mat_f32 +mul_mat_f32_matrix_engine * bisect gemma4: +mul_mat_f16 +mul_mat_f16_matrix_engine * bisect gemma4: +mul_mat_Q8_0 +mul_mat_Q4_0 * bisect gemma4: +flash_attn_ext_f32 +flash_attn_ext_f16_me * bisect gemma4: +mul_mat_id_f32 * bisect gemma4: +sum_rows_f32 * bisect gemma4: +cont_f16 * bisect gemma4: +fill_f32 * bisect gemma4: +unary_f32 (all ops re-enabled except rms_norm_mul) * Update rms_norm_mul_f32.c * bisect2 gemma4 n64: +scale_f32 only * bisect2 gemma4 n64: +rms_norm_f32 +rope_f32 * bisect2 gemma4 n64: +rms_norm_mul_f32 (with ET_UBERKERNEL eviction fix) * bisect2 gemma4 n64: +el_map +get_rows +glu +softmax (skip rms_norm_mul) * bisect2 gemma4 n64: all ops enabled except rms_norm_mul * bisect2 n64: test unary+cont+fill+sum_rows (no mul_mat/flash_attn) * bisect2 n64: +mul_mat_f32 +mul_mat_f32_matrix_engine * bisect2 n64: +mul_mat_f16 +mul_mat_f16_matrix_engine * bisect2 n64: +mul_mat_Q8_0 +mul_mat_Q4_0 * bisect2 n64: +mul_mat_Q8_0 only (disable Q4_0) * bisect2 n64: +mul_mat_Q4_0 only (Q8_0 breaks) * bisect2 n64: +mul_mat_id +flash_attn_ext (skip Q8_0) * run-3: matmul + rms_norm_mul * run-4 * Revert "run-4" * run5 * changes after cleanup * cleanup before upstream * restrict changes into ET backend * move kernel embedding from Python to CMake * move uberkernel gen into CMake * apply clang format * update CMake style * update to match C and C++ style * use source ggml and quant headers instead of ET's * MROPE support * absorb view ops into same branch as none * fix bad rebase * add marty1885 to codeowners * oops * remove redundant newline * fix CI editor warnings --------- Co-authored-by: Vidas <vidas@nuolat.lt> Co-authored-by: Gianluca Guida <glguida@tlbflush.org> Co-authored-by: Gianluca Guida <gianluca@nekko.ai> Co-authored-by: ubergarm <leimgrub@gmail.com> Co-authored-by: SaqibAkram-10xE <saqib.akram@10xengineers.ai> Co-authored-by: Rehan Qasim <rehan.qasim@10xengineers.ai> * hexagon: improve ARGSORT performance for small tensors (llama/25512) * hex-sort: add efficient bitomic sort in hvx regs up to 1024 elements * hex-sort: fix inverted vrors * hex-sort: specialize sort functions for the common cases * hex-sort: add tracing and local context * opencl: add int8 dp4 dense and MoE prefill optimization for Adreno GPUs (llama/25537) * opencl: add int8 dp4 dense and moe GEMM * opencl: refactor --------- Co-authored-by: Li He <lih@qti.qualcomm.com> * Vulkan: route large matmuls to medium tile on Adreno (llama/24877) * [Vulkan] Fixes llama-cli breaking over longer promts sizes The llama-cli was breaking for longer promts sizes for q4_0 quantized networks. Causing due to insufficient shared memory. * Removed the un-used Adreno device * Updated matmul for small pipeline. * ggml : add GGML_OP_LIGHTNING_INDEXER that implements DeepSeek V3.2/V4 lightning indexer (llama/24231) * ggml : add GGML_OP_LIGHTNING_INDEXER that implements DeepSeek V3.2/V4 lightning indexer * ggml : remove scale parameters from lightning indexer OP, add f16 mask parameter * tests : add GGML_OP_LIGHTNING_INDEXER tests * ggml : bump RPC version * chore : check if lightning indexer input tensors are not transposed * tests : count flops instead of bandwidth in lightning indexer test * chore : add missing const * chore : whitespace * ggml : renamed variables in CPU lightning indexer implementation * ggml : fix lightning indexer mask broadcasting * tests : tests for lightning indexer mask broadcasting * chore : whitespace * llama : use GGML_OP_LIGHTNING_INDEXER in DeepSeek V3.2 and DeepSeek V4 models --------- Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com> * cuda: Don't crash when querying memory on device with no free memory. (llama/25157) If a Cuda device has no or limited available memory, the actual call to cudaMemGetInfo() itself can cause a fatal crash due to a cuda out of memory error (there is not enough memory to actually query memory) This causes an issue because we query memory for all devices at startup even if the user isn't trying to use the device for inference. Fix this by making the error non-fatal and assigning zero total/free memory to the device. This will have the downstream effect of the fit algorithm not trying to put any layers on it, which is desired outcome vs hard crashing. this also prevents crashes in cuda enabled builds when user explicitly passes '-dev none' * gguf : reject empty metadata keys (llama/24917) * sycl: add Q2_K to DMMV reorder path (llama/25064) Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * sycl: add fused top-k MoE (llama/25217) * sycl: add fused top-k MoE * sycl: address review: GGML_SYCL_ENABLE_FUSION env, move fusion dispatch to topk-moe * sycl: print GGML_SYCL_ENABLE_FUSION at startup like other env vars Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com> * gguf : add tensor shape accessor (llama/24405) * gguf : add tensor shape accessors * gguf : return tensor shape as const int64_t * * gguf : remove n_dims accessor, keep only gguf_get_tensor_ne * vulkan: Use native e2m1 and e4m3 conversions for mxfp4/nvfp4 (llama/25338) This uses the new VK_EXT_shader_ocp_microscaling_types extension to do fp4 type promotions, and also uses the float8 extension to do ue4m3 promotions for nvfp4. It's reasonable to assume that an implementation that supports fp4 will also support fp8, so we don't need to handle all possible combinations of support. * CUDA: refactor MMQ kernel configuration (llama/24127) * CUDA: refactor MMQ kernel configuration * fix Blackwell config * remove legacy code * metal : add Q2_0 support (llama/25419) * sycl: set fattn_vec_nthreads to 256 for Battlemage (llama/25205) Currently detects lunarlake + battlemage / xe2 and sets the value to 256. Keeps default at 128, Intel's ARC Alchemist's prefered value. * kleidiai : add SME2 f32 kernel (llama/24414) * kleidiai : add SME2 f32 kernel * enable dynamic scheduling for SME2 f32 kernel * ggml: uniformize im2col dst_type for all conv ops (llama/23660) * ggml: uniformize im2col dst_type for all conv ops * Update ggml/src/ggml.c Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * ggml : uniformize im2col casting logic across all conv ops * fix : allow im2col_f16 to accept any kernel type --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * ggml : add a set of functions for checking contiguity of inner tensor dimensions (llama/25650) Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com> * vulkan/cpu: Support f16 as SET_ROWS src. (llama/25432) * vulkan/cpu: Support f16 as SET_ROWS src. This adds full support for f16 SET_ROWS (equivalent to f32) to vulkan and CPU backends, and adds more backend tests. * Set DenormPreserve 16 when supported, to try to fix failures on Intel * tune error threshold * update metal supports_op * opencl: fix a dp4a bug for devices where cl_khr_integer_dot_product is unavailable (llama/25639) * opencl: do not fail backend init on devices without cl_khr_integer_dot_product * opencl: do not call dp4 kernels when dp is unavailable --------- Co-authored-by: Li He <lih@qti.qualcomm.com> * hexagon: fix hmx-queue signal enum-narrowing problem (llama/25677) * opencl: avoid the vec path in GEMV for unaligned row stride (llama/25671) The f16 GEMV kernels take a vectorized path for ne00 >= 128 that casts the row pointers to half4 or float4. When the row stride is not aligned, the wide load becomes misaligned. On devices that require natural alignment for vector loads, the kernel reads garbage. This is the case Intel GPUs and the kernels produce incorrect results there. Adreno happpens to be byte addressable and the kernels happen to work. * opencl: handle OOB write in noshuffle GEMV kernels (odd ne01) (llama/25640) * opencl: do not use `clCreateBufferWithProperties` when targeting CL 2.x (llama/25673) * Flash Attention with XMX engine via oneDNN (llama/25222) * [SYCL] F16 (default) Flash Attention with XMX engine via oneDNN graph API; Qwen3.6-27b-Q8_0 prefill speed up x1.21 at p=512 and x4.26 at p=80k * [SYCL] Address review on FA oneDNN path. Result: llama-bench---pp512; 32% increase with fa1; llama-perplexity---0.11% difference; tested model: mradermacher/Meta-Llama-3.1-8B-Instruct-Q8_0.gguf * PR-25222 revision v2: addressed audits * [SYCL] flash-attn oneDNN SDPA KV F16 rev 3.0: add BMG gate + multi-device sync. Narrow the scrope of this PR to Battlemage only (bmg; Xe2). Other archs (e.g., alchemist) fall back to existing FA kernel. When device_count >1, apply stream -> wait_and_throw(), validated working path for multi-gpu sync fix by @maxious. Co-authored-by: maxious <81432+maxious@users.noreply.github.com> * updated comment on bmg gate, noted the issue --------- Co-authored-by: scientist3 <scientist.3@users.noreply.github.com> Co-authored-by: hmscider <hmscider@users.noreply.github.com> Co-authored-by: maxious <81432+maxious@users.noreply.github.com> * sycl: Increase minimum buffer size for USM system allocations (llama/25525) Raise the threshold for minimum buffer size from 1 GiB to 4 GiB, based on real-world experiments of overcommitting device memory with model weights larger than available VRAM, for example Qwen3.5-35B-A3B-Q8 running on a B70. Also add a debug message to better track USM system allocations. Signed-off-by: Francois Dugast <francois.dugast@intel.com> * sycl : implement xielu op (llama/25550) * sycl : support kernel type fp16 for conv2d_dw (llama/25653) * sycl : fix get_rows Q2_K, Q4_K, Q5_K (llama/25656) * ggml: add f16 out_prod support for CPU and out_prod op for Vulkan (llama/23997) * metal: fuse snake activation (mul, sin, sqr, mul, add) (llama/25459) * metal: fuse snake activation (mul, sin, sqr, mul, add) Mirror the CUDA, Vulkan and CPU snake fusion: same matcher on the naive 5-op chain, same F32 contract on a and inv_b, same F32/F16/BF16 kernel with F32 compute. Follows the Metal backend idioms: bf16 instantiation gated behind GGML_METAL_HAS_BF16 and concurrency ranges checked on the remaining chain nodes before encoding, as done by the bin fusion. Covered by the existing backend-agnostic SNAKE_FUSE tests. * metal: absorb snake fusion into ggml_metal_op_bin Extract the matcher to ggml_metal_op_can_fuse_snake, mirroring the Vulkan naming, and dispatch the fused path from ggml_metal_op_bin. The encode loop switch is back to a single call per case. Address review from ggerganov * metal: fix indentation in ggml_metal_op_can_fuse_snake * cuda : relax tensor contiguity requirements for quantized concat (llama/25678) * cuda : relax tensor contiguity requirements for quantized concat * tests : add test cases for non-contiguous quantized concat * ggml : relax contiguity requirements for quantized concat --------- Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com> * CUDA: tighter MMQ src1 buffer size for native fp4 (llama/25613) * opencl: fix two issues on flash attention for Adreno a7x (llama/25697) * opencl: route `sub_group_shuffle_xor` to qcom ext when KHR ext is unavailable KHR `sub_group_shuffle_xor` is not defined by compiler when `cl_qcom_subgroup_shuffle` is present, causing certain FA kernels fail to build. Define the KHR shuffle_xor using the qcom extension. * opencl: skip FA kernels with mixed and quant types for A7x to avoid compiler crash * cuda : CUDA GGML_OP_LIGHTNING_INDEXER implementation (generic vector kernel + wmma kernel) (llama/25545) * cuda : CUDA GGML_OP_LIGHTNING_INDEXER implementation (generic vector kernel + wmma kernel) * chore : remove indentation of #pragma unroll * cuda : remove unnecessary kernel template declarations * cuda : add WARPS_PER_BLOCK and K_VECS_PER_BLOCK template parameters in lightning indexer kernels to avoid duplication of constants. * cuda : relax MMA architecture requirements to Turing in lightning indexer implementation * chore : renamed variables * chore : rename ggml_cuda_op_lightning_indexer() to ggml_cuda_lightning_indexer() * chore : TODO for AMD rocWMMA * chore : whitespace formatting * chore : another variable rename to fix problems caused by shadowing * chore : yet another rename, this time uppercased all constants * cuda : added alignment checks for Q and K tensors in lightning indexer implementation --------- Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com> * opencl: exclude some moe kernels on Adreno a7x (llama/25698) * opencl: exclude Adreno A7x from using Adreno MoE kernels Some compilers for A7x devices miscompile the repack kernels, corrupting the weights and causing MoE models to generate garbage output * opencl: exclude A6x and unknown Adreno from MoE weights repack * cuda: extract Q1_0 elements via __byte_perm (llama/25628) * opencl: disable FA and MoE weights repack to work around compiler issues for Adreno 850 GPU (llama/25745) * opencl: workaround for A850 compiler compat * opencl: fix DX compiler version parsing and cleanup --------- Co-authored-by: Li He <lih@qti.qualcomm.com> * CUDA: dedup MoE gate/up activation quantization (llama/25441) * CUDA: dedup MoE gate/up activation quantization (fp4) For MoE gate/up projections the src1 activation is broadcast across the routed experts (ne11 == 1), so ids_src1 maps every one of a token's n_expert_used slots to the same physical row. The MMQ path therefore re-quantized each token's activation n_expert_used times. For fp4 (NVFP4/MXFP4) src0, quantize each unique token row once instead of once per expert. For NVFP4 a single quantize+scatter kernel (quantize_scatter_mmq_nvfp4) quantizes each token once and writes the resulting block_fp4_mmq straight to all n_expert_used slots, using an inverse token->compact-row map (build_tok2c). MXFP4, and GGML_CUDA_MOE_QUANT_GATHER=1, use a two-kernel variant: quantize unique rows then gather into the expert-sorted layout (gather_mmq_fp4_blocks). Both are bit-identical to the previous gather-then-quantize path (identical source data, deterministic per-block quantization), verified by test-backend-ops MUL_MAT_ID (type_a=nvfp4, broadcast b=1; 790/790 for the default, gather, and per-expert paths) and by coherent end-to-end generation. Set GGML_CUDA_NO_MOE_QUANT_DEDUP=1 to force the original per-expert path. Same-binary A/B on RTX 5090 (sm_120), Qwen3.6-35B-A3B-NVFP4 prefill @8192 (nsys, graphs-off; the unchanged mul_mat_q GEMM confirms stable clocks): activation-quant GPU-busy drops 61% (78.2 -> 30.4 ms) with the fused quantize+scatter, vs 33% (78.2 -> 52.8 ms) for the two-kernel gather. The fused path avoids materializing and re-reading the 8x compact buffer, writing the expert copies directly from registers. * CUDA: bounds-check token ids in build_tok2c_kernel Guard against malformed ids_src1: skip out-of-range token ids (t < 0 or t >= n_tokens) and drop entries beyond n_expert_used per token instead of writing past the token's tok2c region. No behavior change for valid MoE routing data; test-backend-ops MUL_MAT_ID 790/790. * Refactor the code based on review comments - Removed previously added kernels that were not necessary anymore\ - Added an inverse mapping from (token, slot) to compact row. Each token is quantized once and scattered to its compact rows. * Adding q8_1 support for dedup and addressing review comments * Add pragma unrolls * Remove redundant cudaMemsetAsync call * Removing follow up redundancies --------- Co-authored-by: praneshgo <227579474+praneshgo@users.noreply.github.com> * ggml-cuda : restore prop.integrated on HIP builds (llama/24233) PR #16308 set info.devices[id].integrated = false unconditionally for all CUDA/HIP devices as a workaround for corrupted output on Jetson Orin (#15034). On HIP/ROCm the device's real hipDeviceProp_t.integrated flag is needed: with the cached field forced to false, supports_buft() refuses CUDA host buffers on AMD APU/UMA parts, while get_type() already reads prop.integrated (#23007) — an inconsistency that breaks integrated-GPU host-buffer use on ROCm. Guard the workaround so it only applies to non-HIP (CUDA) builds and restore prop.integrated for HIP, keeping the Jetson workaround intact for CUDA. Fixes #23977 Signed-off-by: liminfei-amd <91481003+liminfei-amd@users.noreply.github.com> * Enable CUDA graphs on volta+turing (llama/25749) * CUDA: Support CUDA Virtual Devices (llama/25228) * support cuda virtual devices * disable NCCL path when virtual devices are used * label virtual devices in description; add GPUx2 server CI jobs * code refactor * vulkan: when using transfer queue for async copies, sync on event_wait to avoid race (llama/25229) * kleidiai: Add SME vs SME2 distinction in kernel dispatch (llama/25478) The current integration treats SME as a single capability (CPU_FEATURE_SME) with no distinction between SME(v1) and SME2. The kernels dispatched under CPU_FEATURE_SME use SME2-specific instructions, making dispatch incorrect on SME(v1)-only hardware. We introduce build-time and runtime distinction between SME and SME2, and wire SME(v1) and SME2 kernels based on actual hardware support. * hexagon: L2 cache handling rework (dirty bit tracking with lazy flushing) and more MUL_MAT updates (llama/25762) * hex-mm: fix artificial limit in the solver that restricted number of act-prep threads * hex-mm: fix warning * hex-prof: do not apply --top to the timeline report * hmx-mm: add suport for tiled act-processing to better distribute hvx work * hex-l2: add tracing for l2flush events * workqueue: redo the legacy workpool api to match hmx-queue and dma-queue * hmx-mm: fix f32 activation buffer alignmnet for nhvx=5,6,7 * hex-work: minor cleanup for work-queue apis * hex-work: further cleanup of the work-queue api * hex-l2: optimize l2flushes at the opbatch level * hex-work: remove unused mask * hex-work: no need to drop hvx ctx in the work-queue * hex-work: add explicit wakeup/suspend and make threads spin * hex-bufs: mark any non-weight tensor as compute * hex-dma: dma-queue support for alias queues and cached dma * hex-l2: track tensor aliases and delay or skip flushes as much as possible * hex-l2: simplify tensor alias handling * hex-l2: handle overlapping views as a circular list of aliases * hex-tens: add flags helper * hex-l2: add helper for marking tensors clearn/dirty * hex-l2: mark binary and rope outputs as l2-clean and keep the rest as is for now * hex-l2: proper support for handling all tensor overlap scenarios * hex-trace: instrument matmul init code and cleanup trace checks * hex-thread: introduce dedicated main thread with explicit stack and priority * hex-l2: track dirty state as bitmap and introduce threaded flush * hex-trace: remove redundant checks for ctx != null * hex-l2: allocate entire context as one buffer and l2fetch it after big flushes * hex-l2: disable tensor clearing in binary and rope for now seems to cause issues with fusion * hmx-mm: update act proc to use fastdivs and fix DMA overflow * hmx-mm: make MUL_MAT_ID kernels robust to multi-chunk cases (start_row>0) * hex-queue: remove obsolete queue interfaces and flush hmx-queue at the end of the op-batch * hex-queue: dont use early wakeup for small op-batches * hex-tensors: properly cap max_tensors in op-batches and dirty_map * hex-l2: make sure threaded l2flush does proper rounding * hex-l2: factor out htp_tensor_flush for reuse (if needed) * hex-l2: optimize tensor flushes by coalescing flush-all * hex-l2: optimize multi-threaded flush * hex-drv: futureproof version checks * hexagon: fix errors and warnings on windows * hex-main: update main thread to only use dspqueue_read, dspqueue_peek is not available on some platforms * hex-main: add fallback mode for dspqueue with callbacks * hex-main: introduce fallback mode for using dspqueue callbacks for full op processing * hex-main: remove early wakeup, not helping and seems to cause some errors with certain batch sizes * hex-l2: make sure to use invalidate version of flushall * hex-l2: dont try to trace early l2flush at the start of op-batch * hex-main: remove offset_ctx that must be zero anyway * hex-hmx: fix hmx_queue_depth to use idx_write - idx_read * hex-hmx: use atomic_load for idx_read/write * hex-main: add static assert to make sure n_threads are aligned * DeepseekV4: Add fused hyper-connection ops (llama/25585) * dsv4 hc-ops * add missing files; * add cparams * update rpc version * address review comments * address review comments * docs: added a note about using OpenCl with Adreno 810 (llama/25786) * opencl: loads quants as uint for q4_K and q5_K flat mv (optimization for Adreno A7x GPUs) (llama/25780) * opencl: load quant as uint in mv_q4_k_f32_flat helps older compilers (e.g., E031.41, boosts 2x), no impact on newer compilers (e.g., E031.45 or newer) * opencl: load quant as uint in mv_q5_K_f32_flat helps older compilers (e.g., E031.41) * opencl: format --------- Co-authored-by: Li He <lih@qti.qualcomm.com> * opencl: add ABS op (llama/25115) * sycl: fix row calculation when K_QUANTS_PER_ITERATION is 1 (llama/25690) * sycl: fix incorrect row calculation when K_QUANTS_PER_ITERATION=1 Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * sycl: use K_QUANTS_PER_ITERATION for non-reordered Q5_K kernel This is the only Q5_K kernel that was not using KQPI. Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * sycl: add missing second half processing to reordered q5_k Error found while running GGML_SYCL_PRIORITIZE_DMMV=1 \ build/bin/test-backend-ops test -o MUL_MAT Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * sycl: fix potential off-by-one error Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * sycl: fix missing row > nrows check Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> --------- Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * vulkan: Support Q2_0 (llama/25430) * vulkan: Support Q2_0 The backend perf tests for mat-vec-mul weren't very good at first (worse than q2_k), doubling the rows per workgroup made a big difference. * reorder * resolve merge conflict, adjust err threshold for f16->q2_0 set_rows * ggml-blas: default hadamard mul_mat to cpu routine (llama/25710) Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * ggml : bump version to 0.17.0 (ggml/1568) * opencl: transpose q4_K noshuffle scales for coalesced reads (llama/25805) * opencl: read/write MoE dp4a activation tiles to local memory as 128-bit (vectorized LD/ST perf opt) for Adreno GPUs (llama/25810) * opencl: read MoE dp4a activation tile as 128-bit local loads * opencl: vectorize MoE dp4a activation staging as 128-bit loads * opencl: load and use `kernel_gemm_moe_q6_k_f32_ns` from bin kernel lib (llama/25797) * opencl: Support broadcast for Adreno MUL_MAT and honor `view_offs` for Adreno Q8_0 MUL_MAT for llama-server multi-stream (llama/25910) * opencl: handle broadcast for adreno gemm/gemv_noshuffle * opencl: honor view_offs for adreno noshuffle gemm/gemv * opencl: general GEMM/GEMV support broadcast * opencl: remove unnecessary tests * opencl: remove unnecessary comments --------- Co-authored-by: Li He <lih@qti.qualcomm.com> * hexagon: add CLAMP op (llama/25934) * CUDA: vectorize same-type get_rows with int4 copy (llama/25929) k_get_rows_float did a scalar one-element-per-thread copy and recomputed the row-invariant work (index load, fast_div_modulo, src/dst row pointers) for every element. Hoist that out of the per-element loop, and add a vectorized path (k_get_rows_float_vec) that copies one int4 (16 B) per thread for the contiguous same-type (no-cast) case. The vectorized path is gated at compile time (is_same<src0_t, dst_t>) and at runtime on 16-byte alignment of the base pointers and all row strides and on ne00 % VEC == 0. Vectorizing divides the block count by VEC, so a small single-row gather can drop below the device CU count and regress; an occupancy gate keeps those on the block-rich scalar path. On Strix Halo (gfx1151) the DeltaNet recurrent-state gather (ne00=524288) drops 18.6us -> 13.0us (rocprofv3 HW timestamps), faster than the Vulkan backend, with no regression on the small conv-state gather; total get_rows -27%. test-backend-ops GET_ROWS passes (47/47). Assisted-by: Claude Opus 4.8 * ggml-openvino: Add GGML_BACKEND_DL_IMPL invocation for OpenVINO backend (llama/25795) This adds the missing `GGML_BACKEND_DL_IMPL()` macro invocation, that other backends have. Fixes #25586 for me * vulkan: Refactor vk_queue to use per-instance mutexes and unique handles (llama/23570) * Refactor vk_queue to use per-instance mutexes and unique handles * integrates VK_KHR_internally_synchronized_queues, abstracting the queue submission into a polymorphic interface that completely bypasses host-side mutex locking when driver-side synchronization is supported * fix compilation error * fix duplicate pNext chain for VkPhysicalDeviceInternallySynchronizedQueuesFeaturesKHR * add fallback defines for VK_KHR_internally_synchronized_queues * add null checks for queues in vk_device_struct destructor * use unique_ptr for outer queues to enforce exclusive ownership and optimize lifetime * use static constexpr for eInternallySynchronizedKHR * add lock guard to ggml_vk_create_aliased_queue for thread safety * initialize sync_query_features.internallySynchronizedQueues to VK_FALSE * reuse sync_query_features for internallySynchronizedQueues and simplify chaining * refactor internallySynchronizedQueues detection * fix internallySynchronizedQueues query guard * use eInternallySynchronizedKHR constant * fix self-referential alias for eInternallySynchronizedKHR * use macro for eInternallySynchronizedKHR fallback * fix internallySynchronizedQueues query timing in ggml-vulkan.cpp to prevent device creation mismatch * reset sync_query_features.pNext before reusing in device creation chain, also removed the redundant second probe call * refactor internally synchronized queues detection to use chained feature query and avoid redundant API calls * Update ggml/src/ggml-vulkan/ggml-vulkan.cpp Co-authored-by: Jeff Bolz <jbolz@nvidia.com> * Update ggml/src/ggml-vulkan/ggml-vulkan.cpp Co-authored-by: Jeff Bolz <jbolz@nvidia.com> * Update ggml/src/ggml-vulkan/ggml-vulkan.cpp Co-authored-by: Jeff Bolz <jbolz@nvidia.com> * Update ggml/src/ggml-vulkan/ggml-vulkan.cpp Co-authored-by: Jeff Bolz <jbolz@nvidia.com> * Update ggml/src/ggml-vulkan/ggml-vulkan.cpp Co-authored-by: Jeff Bolz <jbolz@nvidia.com> * Update ggml/src/ggml-vulkan/ggml-vulkan.cpp Co-authored-by: Jeff Bolz <jbolz@nvidia.com> * rename sync_enable_features to internally_synchronized_queues_features * queue_flags is still computed before has_internally_synchronized_queues is set * fix trailing whitespace * replace eInternallySynchronizedKHR macro with static constexpr * preserve source queue semantics in single-queue aliased transfer queue * vulkan: fix cmd_pool access via pointer for compute_queue unique_ptr * vulkan: lock queue during debug label emission when not internally synchronized --------- Co-authored-by: Jeff Bolz <jbolz@nvidia.com> * kleidiai : warn once when a weight type has no KleidiAI kernel (llama/25701) * cuda: add sqrt_softplus in topk-moe for dsv4 (llama/25896) * hexagon: check tensor type when reusing descriptors (llama/25968) * cuda: GET_ROWS quants (llama/25962) * cuda: add k-quant support to GET_ROWS Device-side embedding lookups require GET_ROWS to handle the k-quants used by common GGUF recipes (Q4_K_M stores token_embd as q6_K). Without it the backend rejects the op and the scheduler falls back to the host, copying the full embedding matrix back on every token in single-device graphs. Factor the super-block dequantizers out of the dequantize_block kernels in convert.cu into shared device functions in dequantize.cuh and reuse them from a new k_get_rows_kq kernel : one thread block dequantizes one (dst row, super-block) pair with the existing thread layouts, 32 threads for q4_K and 64 for the other k-quants. Covers q2_K to q6_K in get_rows_cuda and supports_op. i-quants are left as a TODO. * cuda: add i-quant support to GET_ROWS Extends the shared super-block dequantizers to the nine i-quants and reuses them from k_get_rows_kq with the 32-thread layout of the matching convert.cu kernels. supports_op gates the k-quant and i-quant path on ne0 being a multiple of QK_K, which iq4_nl does not guarantee on its own (QK4_NL sub-blocks). mxfp4 is left as a TODO. * cuda: add mxfp4 support to GET_ROWS Moves the mxfp4 dequantizer into the shared super-block helpers and reuses it from k_get_rows_kq with the 32-thread layout of the matching convert.cu kernel. mxfp4 joins the ne0 % QK_K gate in supports_op since its 32-value sub-blocks do not guarantee QK_K-aligned rows on their own. This closes GET_ROWS type coverage on CUDA: every quantized GGML type now takes the direct device path. * cuda: gate the GET_ROWS row size only for 32-value sub-block types Address review from @pwilkin: the i-quant commit replaced the return shared by the whole supported type cascade, so f16/f32/bf16/i32 and the legacy quants also inherited the ne0 % QK_K == 0 gate and any row size that is not a multiple of 256 fell back to the scheduler. Split the cascade: unconditional support is restored everywhere, the gate stays only on iq4_nl and mxfp4 whose 32-value sub-blocks do not guarantee the QK_K super-blocks the kernel iterates on. * webgpu : add CONV_2D_DW (depthwise conv2d) kernel (llama/25847) * webgpu : add CONV_2D_DW (depthwise conv2d) kernel Implement GGML_OP_CONV_2D_DW for the WebGPU backend, ported from the Vulkan backend's conv2d_dw.comp. Assisted-by: Claude Opus-4.8 * Remove unnecessary comments in webgpu support * update supported ops tables, triggered by adding webgpu CONV_2D_DW * ggml: enable PowerPC backend variants on AIX (llama/25983) * ggml: enable PowerPC backend variants on AIX Allow the PowerPC CPU backend variants to be built on AIX by extending the platform check in the CMake configuration. This reuses the existing PowerPC backend implementations without changing their behavior. Also fix a missing semicolon in the PowerPC Q0 matmul implementation. * Fix missing semicolon in sgemm.cpp * hexagon: activation ops update (llama/25974) * hex-geglu: optimized all-in-one geglu microkernel * hex-geglu: enable non-contiguous src and strided DMA * hex-act: enable non-contiguous srs and strided DMA for rest of ACT ops * hex-act: generalize GLU per-thread functions via DEFINE_GLU_PER_THREAD macro * hexagon: move UNARY_SILU and UNARY_GELU to unary-ops * hex-act: replace the generic ops_context scratchpad usage with a local htp_vtcm_layout computation per act op. --------- Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com> * CUDA: Improve NVFP4 W4A4 activation quantization (llama/25730) * Squash history before conflict-resolution during rebase on master WIP commit Add 32-byte loads, restore per-block amax Use nvfp4x4 intrinsic when available Fuse per-channel amax and quantization kernels Do pointer arithmetic only once on x Remove unnecessary ternary in the load We assert on host side that ne00 is 64-aligned Add back scale-search, but optimize it with intrinsics Code cleanup Make scale in MMQ-epilogue NVFP4-specific/restrictive for now Remove unneeded include, add comment Fix trailing whitespace Guard __builtin_align__(32) struct to NVIDIA Seems like HIP doesn't have this available, see https://github.com/ggml-org/llama.cpp/actions/runs/29438651734/job/87431623001 * compiler massaging to avoid unnecessary LDCs * kvalues_mxfp4 -> kvalues_nvfp4 in quantize_mmq_nvfp4 * Always pass in src1_scale.ptr * Extract ggml_cuda_is_aligned helper * metal : add f16 type support to leaky relu (llama/25981) * CUDA: fix external compilation of q1_0 MMQ (llama/25778) * hexagon: fix Windows crash when op_poll is enabled (llama/26029) * hexagon: further improved pipeline of the core bits (L2, DMA, MM, FA) (llama/26049) * hex-l2: use dirty ranges for flushing * hex-l2: simplify range based flush logic * hex-l2: optimize dirty range scans * hex-hvx: support for reduce_max_i32 * hex-mm: optimize fused MUL_MAT+ADD to use vtcm for bias when it fits * hex-mmid: optimize mmid row-mapping generation * hex-mmid: optimize mmid row-mapping generation * hex-mmid: optimize mmid row-mapping generation (round2) * hmx-mm: optimize output proc by tiling (col-chunking) * hex-fa: start the next q dmas a bit earlier * hex-fa: prefetch Q even earlier * hvx-fa: optimize softmax to keep things in hvx registers * hex-fa: hoist const register init in softmax loop * hmx-fa: kick off next-qkv DMAs before o-proc * hmx-fa: hoist various checks out of the inner loop * hmx-fa: adjust the cost model to better balance softmax work across hvx threads * hmx-fa: overlap diag rescale build with last HMX task * hmx-fa: optimize idx update in output proc * hmx-fa: unroll the softmax loops for improved perf * hmx-fa: overlap qk-dot with softmax, double-buffer p and s tiles * hex-trace: double the default number of trace entries * hex-trace: add trace events for opbatch and buffer mgmt * hex-trace: overhaul tracing to simplify runtime event handling and support opbatch stats * hex-trace: replace ascii timeline diagram with pipeline bubbles detector * hex-trace: handle missing start/stop events * hex-dma: always log stop/start trace events even for dummy dmas * hex-scripts: fix flake warnings * opencl: do not treat NULL-mask flash attention as causal (llama/25771) * opencl: cache compiled cl_program binaries on disk (llama/26050) * HIP: remove rocWMMA FlashAttention (llama/26046) * Update ggml/src/gguf.cpp : Defined virtual keyword for destructor of gguf_writer_base (llama/25867) Without a virtual destructor, deleting a derived object through a base-class pointer only invokes the base destructor, skipping the derived one. * hexagon: partial im2col support (llama/26007) * hexagon: add IM2COL op Add Hexagon IM2COL support targeting only patch-embedding convolutions. * hexagon: im2col refactor and cleanup * hex-im2col: instrument and update im2col. * hex-im2col: add local htp_vtcm_layout computation. * opencl: fix fused RMS norm mul view offset (llama/26085) * ggml-cpu: Enable BF16 tiled gemm optimization on PowerPC (llama/26068) * ggml : adjust logic for offloading ops to weight's backend (llama/25832) * ggml : adjust logic for offloading ops to weight's backend * llama : dsv4 graph fixes * sycl(build): parallelize ocloc invocations (llama/25903) * Disable -ffast-math on HIP (llama/25495) * ggml-metal: FWHT kernel for metal backend (llama/25924) * metal fwht wip * shape guard and formatting * formatting * Formatting and typos Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com> * fix narrowing issue Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com> * cont : minor style --------- Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * sycl: fix use-after-return of the SDPA scale in the oneDNN flash-attention path (llama/25880) * sycl: fix use-after-return of the SDPA scale in the oneDNN flash-attention path The scale was uploaded with an async memcpy sourced from a stack local. On the in-order queue that copy is ordered behind the K/V staging kernels; once n_kv is large enough (>= ~26k observed on Arc Pro B70) the staging outlives the host stack frame and the copy reads recycled memory, feeding the SDPA a garbage scale. Output then collapses to a single repeated token and the KV cache is poisoned for the rest of the session. Short contexts win the race by accident, and test-backend-ops caps FLASH_ATTN_EXT at kv=1024, which is why CI never caught it. The previous device_count > 1 wait_and_throw() gate (and reverting it, PR #25741) fixes the symptom only by keeping the frame alive across the copy at the cost of a host sync on every FA call. Fix: cache one device scalar per (device, value) -- the scale is constant per model -- and upload it synchronously once. The single-device fast path (no per-call host sync) is then safe: every device-side hazard already serializes on the in-order queue. The multi-GPU conservative wait is kept unchanged. Also: - GGML_SYCL_FA_ONEDNN_MAX_KV env (0 = unlimited): optional n_kv ceiling that routes very long sequences to the native FA kernel. - test-backend-ops: FLASH_ATTN_EXT F16 cases up to kv=65536 (Qwen3.6-27B geometry hsk=hsv=256 GQA 6, and hsk=128 GQA 4), closing the kv=1024 blind spot. Note the race itself needs a live multi-op pipeline to reproduce; single-op runs pass even on broken builds. Verified on Arc Pro B70 (bmg_g31), Qwen3.6-27B Q4_K, -c 131072: output byte-identical at temp 0 to the native FA path through 32k-deep prefill, with prefill depth-flat at 820-840 t/s (vs 340-350 native at 32k depth). Assisted-by: Claude Fable 5 * sycl: handle GGML_SYCL_FA_ONEDNN_MAX_KV like the other runtime env vars and document it Review feedback on #25880: - read the variable once at backend init into g_ggml_sycl_fa_onednn_max_kv via ggml_sycl_get_env, and print it in the startup env listing (-lv 4 shows it) - document GGML_SYCL_FA_ONEDNN and GGML_SYCL_FA_ONEDNN_MAX_KV in the SYCL.md runtime table Also trim the added FLASH_ATTN_EXT cases to kv={4096,16384}: the 32768/65536 shapes exceed the legacy NMSE threshold on both the oneDNN and native kernels (long-sequence fp16 accumulation drift, present before this PR) and would fail CI for an unrelated reason. Assisted-by: Claude Fable 5 * sycl: clarify GGML_SYCL_FA_ONEDNN_MAX_KV default is disabled Assisted-by: Claude Fable 5 * sycl: state default behavior of GGML_SYCL_FA_ONEDNN_MAX_KV explicitly Assisted-by: Claude Fable 5 * Update ggml/src/ggml-sycl/fattn-onednn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * sycl: write the SDPA scale from a kernel instead of caching it The per-(device, value) scale cache was a function-local static unordered_map with no synchronization, so concurrent backend instances could access and rehash it at the same time. Write the scalar with a single_task instead. The value is captured into the command, so no host memory has to outlive the call -- which is what the use-after-return fix needed in the first place. That removes the shared container, the leaked device allocation and the string key, and it also closes the remaining async-memcpy-from-a-stack-local on the first flash-attention call. Ordering does not rely on timing: the queue is created with sycl::property::queue::in_order and the dnnl stream wraps that same queue, so the write completes before the SDPA reads the scalar. The multi-GPU wait_and_throw() branch is unchanged. Also drop the <cstdlib> include, which is unused. Assisted-by: Claude Opus 5 --------- Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * ggml-cuda: add chunked SSD matmul for Mamba-2 prefill acceleration (llama/22675) * ggml-cuda: add chunked SSD matmul for Mamba-2 prefill acceleration * cuda: added SSD CICD fixes for CUDA / HIP / MUSA / MSVC. * ggml-cuda: review comments fixed. * ggml-cuda: Fuse M matrix materialization into pre_matmul kernel and enabled test. * ggml-cuda: test updates and fixes * ggml-cuda: test updates to remove hardcoding of tensor initialise data limits. * ggml-cuda: ssd minor review comment fixed. * ggml-cuda: ssd minor CICD fixed. * CUDA SSD: Fixes correctness by promoting s0_stride_seq to int64_t, improves memory coalescing in ssm_ssd_prepare_dt_kernel, and boosts efficiency by merging B_weighted and C_scaled; also addresses prior review comments. * cuda: fix sdata read-write race in prepare_dt fallback scan loop * vulkan: add iq4_nl support back to FA (llama/24585) * vulkan: add iq4_nl support back to FA I was originally concerned about wasting shared memory on the LUT, but it's small and unlikely to matter in practice. Also support q1_0 for non-coopmat2. Fixes #23681 * remove q1_0 FA support * ggml : set output of view src (llama/25729) * llama-graph: set_outputs to t->view_src * change set_output to GGML_ASSERT about views not being outputs * sampler : avoid views in outputs * cont : fix dist sampler * cont : consistent logits handling * ggml : set output of view src * graph : simplify set_outputs() * cont : cleanup Co-authored-by: Gaurav Garg <gaugarg@nvidia.com> --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> Co-authored-by: Gaurav Garg <gaugarg@nvidia.com> * opencl: skip the Adreno KQ/KQV image kernels for multi-stream batches (llama/26189) The Adreno KQ/KQV image1d kernels (ggml_cl_mul_mat_kq_kqv_adreno) ignore dim 3 entirely: the sub-buffer covers only nb02*ne02 bytes and the kernel receives no ne03/ne13/nb03/nb13 arguments. With the unified KV cache, multi-sequence batches (e.g. llama-perplexity with its default -b 2048, n_seq=4, or a multi-slot llama-server) present KQ/KQV as 4D tensors with ne3 = n_stream, so every stream past the first reads the first stream's K/V and produces garbage. Flash attention masks the bug where it is enabled; devices where FA is declined (e.g. Adreno 740) hit it with default settings. Route ne03/ne13 > 1 to the general path, which handles dim 3, and honor view_offs when creating the sub-buffers (currently always 0 for tensors reaching this function, but the function would silently misread any future view). Llama-3.2-1B-Instruct Q4_0, wiki.test.raw, 8 chunks, -ngl 99: - Adreno 740, default: PPL 1817.64 -> 15.61 - Adreno 740, -fa 0: PPL 1941.64 -> 15.61 - Adreno 840, -fa 0: PPL 1943.90 -> 15.50 - single-stream (-b 512) results unchanged (15.6090) - test-backend-ops -o MUL_MAT on 740: identical before/after (909 OK, 12 pre-existing q6_K failures) * ggml-webgpu: Fix some binding alias issues to support all archs, fix recurrent-state-rollback test (llama/25931) * Add overlap glu variant to support all archs, fix recurrent-state-rollback test * format * Fix all arch overlapped ranges * format * diagnose bus error on apple ci * More testing * more testing * more targeted testing * Fix bug in alignment for > 4gb buffer offsets * Fix bug in view offsets * Try avoiding multi_buffers * not fixed yet, more logging :( * Handle edge case in set_rows * Try looking at view source * Skip deepseek32 for now and clean up trace infrastructure * simplify skipping * last cleanup * actually final cleanup * update handling of overlap * format * try skipping other failing model * add rdna3.5, and 3 to mmq configs so they can be tuned independently. (llama/26199) * RPC: add tensor_memset (llama/25912) * sycl: contiguous fast path + 32-bit index math for unary elementwise ops (llama/25946) * sycl: contiguous fast path + 32-bit index math for unary elementwise ops * sycl: use fastdiv for elementwise index math * ggml-cuda : disable MMQ on devices with less than 48 KiB shared memory (llama/26141) ggml_cuda_should_use_mmq() selects MMQ purely from the quantization type. The current MMQ configurations are designed and maintained against a minimum of 48 KiB per-block shared memory, the limit provided by NVIDIA Pascal GPUs and later. On devices that report less, no supported MMQ tile fits and mul_mat_q_switch_J() aborts when every tile size exceeds the device's per-block shared memory budget. Disable MMQ when smpbo < 48 KiB so the caller falls back to the BLAS path instead of hitting GGML_ABORT. Some current MUSA QY1 devices report only 28 KiB and are covered by this guard. Reproduced on a Moore Threads MTT S70 (arch mp_21, 28 KiB shared memory per block) with an RWKV-7 0.1B Q8_0 model: $ llama-bench -m rwkv7-g1d-0.1b-Q8_0.gguf -p 128 -n 0 J_best=0 ggml/src/ggml-cuda/template-instances/../mmq.cuh:1521: fatal error (core dumped) Only prefill (batch > 1) is affected; token generation is fine. After the fix the same device falls back to the BLAS path: Q8_0 pp128 1470.7 t/s, tg8 55.3 t/s (was: abort) FP16 unchanged Q4_K_M unchanged This matches a -DGGML_CUDA_FORCE_CUBLAS=ON build (pp128 1464.2 t/s), which confirms the fallback path is the one being taken. This is not MUSA-specific: any device with less than 48 KiB per-block shared memory is affected. Co-authored-by: KakaruHayate <KakaruHayate@users.noreply.github.com> * ggml : Fix issue with kleidiai ci and stringop overflow warning (llama/26277) Signed-off-by: Jonathan Clohessy <Jonathan.Clohessy@arm.com> * metal: fix memory unwire if model is freed without any GPU operations (llama/26082) * metal: fix memory leak if model is freed without any GPU operations * metal: run dummy work only if residency sets are used * metal: wrap function in #if defined * metal: measure system-wide wired memory in test * metal: always build regression test Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com> --------- Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com> * CUDA: add Q2_0 support (llama/25707) * ggml : bump version to 0.18.0 (ggml/1576) * sync : ggml * parakeet : update parakeet test models generation This commit updates the parakeet test models to ensure that the random values generated for mel filters are not negative. This will otherwise cause a debug assertion and fail the parakeet-test. Refs: https://github.com/ggml-org/whisper.cpp/actions/runs/30547589734/job/90887624319?pr=3962 --------- Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> Signed-off-by: Francois Dugast <francois.dugast@intel.com> Signed-off-by: liminfei-amd <91481003+liminfei-amd@users.noreply.github.com> Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> Signed-off-by: Jonathan Clohessy <Jonathan.Clohessy@arm.com> Co-authored-by: Aparna M P <aparmp@qti.qualcomm.com> Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com> Co-authored-by: fairydreaming <166155368+fairydreaming@users.noreply.github.com> Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com> Co-authored-by: Johannes Gäßler <johannesg@5d6.de> Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com> Co-authored-by: Martin Chang <marty1885@users.noreply.github.com> Co-authored-by: Vidas <vidas@nuolat.lt> Co-authored-by: Gianluca Guida <glguida@tlbflush.org> Co-authored-by: Gianluca Guida <gianluca@nekko.ai> Co-authored-by: ubergarm <leimgrub@gmail.com> Co-authored-by: SaqibAkram-10xE <saqib.akram@10xengineers.ai> Co-authored-by: Rehan Qasim <rehan.qasim@10xengineers.ai> Co-authored-by: Li He <lih@qti.qualcomm.com> Co-authored-by: Raman Shinde <raman.shinde15@gmail.com> Co-authored-by: cphlipot <9103367+cphlipot@users.noreply.github.com> Co-authored-by: Rohit Mahesh <74331568+rohitmahesh1@users.noreply.github.com> Co-authored-by: Todd Malsbary <todd.malsbary@intel.com> Co-authored-by: Frosty40 <newjordan@gmail.com> Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Co-authored-by: QuintinShaw <yx6f20@soton.ac.uk> Co-authored-by: Jeff Bolz <jbolz@nvidia.com> Co-authored-by: Pasha Khosravi <khosravipasha@users.noreply.github.com> Co-authored-by: Titaniumtown <titaniumtown@proton.me> Co-authored-by: Charles Xu <charles.xu@arm.com> Co-authored-by: JusteLeo <leonard.adamo66@gmail.com> Co-authored-by: Chyan <163109379+chyan8@users.noreply.github.com> Co-authored-by: hmscider <201289679+hmscider@users.noreply.github.com> Co-authored-by: scientist3 <scientist.3@users.noreply.github.com> Co-authored-by: hmscider <hmscider@users.noreply.github.com> Co-authored-by: maxious <81432+maxious@users.noreply.github.com> Co-authored-by: Francois Dugast <francois.dugast@intel.com> Co-authored-by: Andrew Smith <atsmith19@comcast.net> Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> Co-authored-by: Michael Lamothe <michael.lamothe@gmail.com> Co-authored-by: Pascal <admin@serveurperso.com> Co-authored-by: leonardHONG <2695316095@qq.com> Co-authored-by: David Friehs <david@friehs.info> Co-authored-by: Pranesh Gonegandla <pranesh.iitp@gmail.com> Co-authored-by: praneshgo <227579474+praneshgo@users.noreply.github.com> Co-authored-by: liminfei-amd <91481003+liminfei-amd@users.noreply.github.com> Co-authored-by: Alexander Heisler <126129661+heislera763@users.noreply.github.com> Co-authored-by: Anav Prasad <anavp@nvidia.com> Co-authored-by: Ruben Ortlam <rortlam@redhat.com> Co-authored-by: Rajendra Matcha <matcraje@qti.qualcomm.com> Co-authored-by: Aman Gupta <amangupta052@gmail.com> Co-authored-by: akleine <alb.kleine@gmx.de> Co-authored-by: Gezahegne <gezahegne.yirefu@gmail.com> Co-authored-by: Aaron Teo <aaron.teo1@ibm.com> Co-authored-by: Todor Boinovski <todorb@qti.qualcomm.com> Co-authored-by: Piotr Wilkin (ilintar) <piotr.wilkin@syndatis.com> Co-authored-by: Markus Ebner <seijikun@users.noreply.github.com> Co-authored-by: Winston Ma <winstonma@ymail.com> Co-authored-by: Kamalesh VS <76260512+kkjjkamal123@users.noreply.github.com> Co-authored-by: Wei Wang <w10493wang@163.com> Co-authored-by: m1el <m1el@ya.ru> Co-authored-by: shalinib-ibm <Shalini.Salomi.Bodapati@ibm.com> Co-authored-by: Oliver Simons <osimons@nvidia.com> Co-authored-by: Ilia Ilmer <iliailmer@users.noreply.github.com> Co-authored-by: adgup-qti <adgup@qti.qualcomm.com> Co-authored-by: kumaal <44551860+kumaal@users.noreply.github.com> Co-authored-by: Yongmin Yoo 유용민 <yymin1022@gmail.com> Co-authored-by: yzyyzyhhh <96101183+happyyzy@users.noreply.github.com> Co-authored-by: Beinsezii <39478211+Beinsezii@users.noreply.github.com> Co-authored-by: Nick Lafleur <55208706+nicklafleur@users.noreply.github.com> Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com> Co-authored-by: meatposes <computerdork@verizon.net> Co-authored-by: Bhavik Sharda <10757940+BLSharda@users.noreply.github.com> Co-authored-by: Gaurav Garg <gaugarg@nvidia.com> Co-authored-by: Reese Levine <reeselevine1@gmail.com> Co-authored-by: Geramy Loveless <gloveless@jqluv.com> Co-authored-by: Kakaru <97896816+KakaruHayate@users.noreply.github.com> Co-authored-by: KakaruHayate <KakaruHayate@users.noreply.github.com> Co-authored-by: Jonathan Clohessy <jonathan.clohessy@arm.com> Co-authored-by: Niklas Wenzel <dev@nikwen.de> Co-authored-by: Daniel Bevenius <daniel.bevenius@gmail.com>	2026-07-31 09:11:28 +02:00
Pasha KhosraviandGeorgi Gerganov	66e178ec84	Add Q2_0 quantization: type definition and CPU backend (llama/24448)	2026-07-10 13:06:42 +03:00
Gaurav GargandGeorgi Gerganov	1a1900f90c	Remove padding and multiple D2D copies for MTP (llama/24086) ... * Make ggml_gated_delta_net take only the initial recurrent state (D, 1, n_seqs) and passes the snapshot count K as an op parameter instead of inferring it from state->ne[1]. Remove the padding hack and copy all emitted snapshots into the recurrent cache with a single strided ggml_cpy * Make GDN changes in all backends. Address review comments. * Fix CI build errors	2026-06-15 10:33:53 +03:00
PascalandGeorgi Gerganov	28c7ed3db7	ggml : add GGML_OP_COL2IM_1D (llama/24206) ... * cpu: add GGML_OP_COL2IM_1D Add the overlap-add (scatter-add) step of a 1D transposed convolution. A ConvTranspose1d factorizes as a GEMM followed by col2im: a weight pre-permuted to [IC, K*OC] is contracted against the [IC, T_in] input with mul_mat to produce a column matrix [K*OC, T_in], and col2im_1d scatters those columns back into the [T_out, OC] signal, with T_out = (T_in - 1)*s0 + K - 2*p0. Keeping the contraction as a plain mul_mat leaves the heavy work on the optimized (and quantizable) matmul kernels, so col2im_1d only does the cheap overlap-add. CPU uses a gather formulation parallelized over output channels, supporting F32, F16 and BF16 with an F32 accumulator. * tests: add backend coverage for GGML_OP_COL2IM_1D Add test_col2im_1d next to the conv_transpose_1d cases, covering F32, F16 and BF16 across eight geometries: the canonical kernel = 2*stride DAC upsampling shape, overlap, no overlap, cropping (p0 = 1 and p0 = stride/2), kernel < stride with zeroed gaps, kernel not a multiple of stride, and a single column unfold. Perf mode gets three real vocoder stage shapes reporting memory bandwidth. max_nmse_err relaxes to 5e-4 for F16 and BF16. * cpu: harden GGML_OP_COL2IM_1D ggml_col2im_1d validates s0, oc, p0 and input contiguity at graph build time, before the oc division, protecting every backend at once. The kernel asserts the contiguity its flat indexing assumes and its doc states the full output length including the crop term. The kernel parallelizes over the time axis: the split stays balanced down to OC = 1, where the previous channel split was single threaded. Values are bit identical on the three real vocoder chains, two out of three improve. * tests: extend the GGML_OP_COL2IM_1D grid The eval grid grows to eleven geometries: OC = 1 (mono output stage), K = 1 with stride > 1 (sparse scatter, every gap position zeroed) and a crop down to T_out = 2 where all the gather bounds act at once. * tests: add col2im_1d equivalence test tests/test-col2im-1d.cpp proves mul_mat + col2im_1d matches the native ggml_conv_transpose_1d on the CPU backend, F32 bit exact, F16 and BF16 through casts of the column matrix. test-backend-ops cannot cover this for a CPU only op since the CPU backend is its own reference there. * rpc: bump protocol patch version for GGML_OP_COL2IM_1D GGML_OP_COUNT goes from 96 to 97 with the new op, which trips the static_assert in ggml-rpc.h. Bump RPC_PROTO_PATCH_VERSION since the op is appended and no existing op code shifts.	2026-06-15 10:33:53 +03:00
Ori PekelmanandGeorgi Gerganov	e78e693017	ggml.h: correct ggml_silu_back arg docstring (a=dy, b=x) (ggml/1500)	2026-05-25 12:26:07 +03:00
Aman GuptaandGeorgi Gerganov	23f956de33	llama + spec: MTP Support (llama/22673) ... * spec: support MTP * fix batch size * rename files * cont : simplify (llama/7) * MTP: clean-up (llama/9) * MTP: clean-up * review: use llama_context_type instead of llama_graph_type * review: remove llama_model_has_mtp * review: fix convert issues * convert: fix pycheck * review: formatting * use `mtp-` for identifying mtp models * convert: fix mtp conversion * mtp -> draft-mtp * remove unused llama_arch * add need_embd in speculative * llama: allow partial seq_rm for GDN models for speculative decoding Currently speculative checkpoint needs to restart from a checkpoint after some draft tokens are not accepted, this leads to some wastage in running the target again. This PR adds the ability to rollback upto `draft_max` by storing the GDN intermediates. * fix pending state * vulkan: add GDN partial rollback * meta: extend check to axis 1 * metal: add GDN partial rollback Extend the gated delta net kernel to store intermediate states for partial rollback support on the Metal backend. - Add K (snapshot slot count) as a function constant - Read input state from slot 0 of the 3D state tensor - Write intermediate states to different slots during token loop - For K=1, maintain backward-compatible single-slot behavior Ref: https://github.com/ggml-org/llama.cpp/commit/8c05923630110223669f069af2000e9cf10c02bc Assisted-by: llama.cpp:local pi * delta_net_base: use ggml_pad instead of new_tensor * review: add need_rs_seq * review: rename part_bounded to n_rs * review: deslop comments * review: rename, add asserts * server : adjust checkpoint logic (llama/11) * server : adjust checkpoint logic * cont : rm asserts * server-context: fix early exit * spec : fix compatibility with n-gram and add TODOs (llama/13) * metal : cleanup * llama : fix faulty bitwise check in recurrent memory * server : disable RS-based MTP in combination with other spec types * spec : add TODOs * cont : fix comment * cont : update comment * common : fix logic for ngram + mtp compat * llama-memory: enable checkpointing with partial rollback * cont: add test-case for loading into a dirty ctx * llama-memory-recurrent: clear rs_idx in clear * download: fix mtp path * llama-arch: fix enorm op * docs: update docs * conversion: fix type annotations --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>	2026-05-25 12:26:07 +03:00
IsmailandGeorgi Gerganov	4794432337	ggml : implement fast walsh-hadamard transform for kv rotation (#21352) (llama/22631)	2026-05-14 21:26:48 +03:00
Xuan-Son NguyenandGeorgi Gerganov	86d94cd95b	docs: more extensive RoPE documentation [no ci] (llama/21953) ... * more extensive ggml_rope documentation * add more docs * nits	2026-04-30 11:29:09 +03:00
Georgi Gerganov	f1d2b83db0	ggml : deprecate GGML_OP_ADD1 (llama/21363) ... * ggml : deprecate GGML_OP_ADD1 * cont : remove tests * cont : re-enable vulkan check	2026-04-30 11:29:02 +03:00
3 people	0c2fbd4703	ggml: add Q1_0 1-bit quantization support (CPU) (llama/21273) ... * ggml: add Q1_0 and Q1_0_g128 1-bit quantization support (CPU) * add generic fallback for x86 * remove Q1_0 (group size 32) * rename Q1_0_g128 => Q1_0 * fix Q1_0 LlamaFileType Enum * Fix trailing spaces; add generic fallback for othre backends * Apply suggestions from code review Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com> * fix /r/n spacing + arch-fallback --------- Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>	2026-04-30 11:29:01 +03:00
Georgi Gerganov	945d3151d9	ggml : restore ggml_type_sizef() to aboid major version bump (ggml/1441)	2026-03-18 15:18:24 +02:00
6 people	d73fe25267	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 <amangupta052@gmail.com> * cont : fix the fix Co-authored-by: Aman Gupta <amangupta052@gmail.com> * 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 <noreply@anthropic.com> * metal : validate contiguity of all input tensors in supports_op Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * metal : add algorithm equivalence comment for GDA decay path Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * cont : unslop + optimize * cont : clean-up --------- Co-authored-by: Paul Flynn <paul@arkavo.com> Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> * 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 <devnull@uvos.xyz> * llama : refactor llm_build_delta_net_base API --------- Co-authored-by: Aman Gupta <amangupta052@gmail.com> Co-authored-by: Paul Flynn <paul@arkavo.com> Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> Co-authored-by: Oliver Simons <osimons@nvidia.com> Co-authored-by: uvos <devnull@uvos.xyz>	2026-03-16 13:10:15 +02:00
3 people	e4021d4071	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 <sigbjorn.skjaeret@scala.com> * 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 <sigbjorn.skjaeret@scala.com> * Use fallback quant config * Simplify Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com> * organize * Refactor * Update convert_hf_to_gguf.py Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com> * Update convert_hf_to_gguf.py Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com> * Update convert_hf_to_gguf.py Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com> * 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 <sigbjorn.skjaeret@scala.com>	2026-03-16 13:10:15 +02:00
Aman GuptaandGeorgi Gerganov	49489bfbd1	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 <ggerganov@gmail.com>	2026-03-16 13:10:15 +02:00
4 people	67abc63e9d	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 <mail@marcelpetrick.it> * Update docs/backend/CANN.md Co-authored-by: Aaron Teo <taronaeo@gmail.com> * 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 <sigbjorn.skjaeret@scala.com> * Update tests/test-backend-ops.cpp Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com> --------- Signed-off-by: Marcel Petrick <mail@marcelpetrick.it> Co-authored-by: Aaron Teo <taronaeo@gmail.com> Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>	2026-03-16 13:10:15 +02:00
Georgi GerganovandSigbjørn Skjæret	279be33a83	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 <sigbjorn.skjaeret@scala.com>	2026-02-27 20:57:58 +02:00
JuddandGeorgi Gerganov	5ee5748722	ggml : make ggml_is_view as API (llama/19539) ... * make `ggml_is_view` as API * introduce `ggml_aux_is_view` as inline version for internal use. * change `ggml_aux_is_view` to `ggml_impl_is_view`	2026-02-27 20:57:58 +02:00
Christian KastnerandGeorgi Gerganov	0e219ebf89	docs : Minor cleanups (llama/19252) ... * Update old URLs to github.com/ggml-org/ * Bump copyrights	2026-02-08 09:29:10 +02:00
Georgi Gerganov	47f3e3b927	ggml : add ggml_build_forward_select (llama/18550) ... * ggml : add ggml_build_forward_select * cuda : adapt CUDA graph compat to new feature * vulkan : update logic to handle command buffer closing * ggml : check compute for fusion * ggml : add comment	2026-01-30 15:56:40 +02:00
Masashi YoshimuraandGeorgi Gerganov	a71127dfd8	ggml-webgpu: Fix GGML_MEM_ALIGN to 8 for emscripten. (llama/18628) ... * Fix GGML_MEM_ALIGN to 8 for emscripten. * Add a comment explaining the need for GGML_MEM_ALIGN == 8 in 64-bit wasm with emscripten	2026-01-14 09:11:59 +02:00
Johannes GäßlerandGeorgi Gerganov	aaf3f39b4a	llama: automatically set parameters not set by the user in such a way that maximizes GPU utilization (llama/16653) ... * llama: automatically fit args to free memory llama-fit-params tool * fix CI * hints for bug reports, ensure no reallocation * fix segfault with Vulkan * add llama-fit-params to CI * fix CI * fix CI * fix CI * minor adjustments * fix assignment of 1 dense layer * fix logger not being reset on model load failure * remove --n-gpu-layer hint on model load failure * fix llama-fit-params verbosity * fix edge case * fix typo [no ci]	2025-12-18 08:20:56 +02:00
Georgi Gerganov	cd9b8c6d18	ggml : remove GGML_KQ_MASK_PAD constant (llama/17910) ... * ggml : remove GGML_KQ_MASK_PAD constant * cont : remove comment	2025-12-12 17:53:24 +02:00
5 people	c5e1807071	ggml : add circular tiling support to pad, for Vulkan, CUDA, and CPU (used for making seamless textures) (llama/16985) ... * Feat: Added vulkan circular tiling support * Feat: Added cpu circular * Feat: Added cuda kernels * Added tests * Added tests * Removed non-pad operations * Removed unneded changes * removed backend non pad tests * Update test-backend-ops.cpp * Fixed comment on pad test * removed trailing whitespace * Removed unneded test in test-backend-ops * Removed removed test from calls * Update ggml/src/ggml-vulkan/vulkan-shaders/pad.comp Co-authored-by: Ruben Ortlam <picard12@live.de> * Fixed alignment * Formatting Co-authored-by: Aman Gupta <amangupta052@gmail.com> * Format pad * Format * Clang format * format * format * don't change so much stuff * clang format and update to bool * fix duplicates * don't need to fix the padding * make circular bool * duplicate again * rename vulkan to wrap around * Don't need indent * moved to const expr * removed unneded extra line break * More readable method calls * Minor wording changes * Added final newline * Update ggml/include/ggml.h Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Update ggml/include/ggml.h Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Added circular pad ext tests * Gate non circular pad devices * Cleaned gating of non-circular pad devices --------- Co-authored-by: Phylliida <phylliidadev@gmail.com> Co-authored-by: Ruben Ortlam <picard12@live.de> Co-authored-by: Aman Gupta <amangupta052@gmail.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>	2025-12-12 17:53:20 +02:00
Adrien GallouëtandGeorgi Gerganov	92e50155c9	build : move _WIN32_WINNT definition to headers (llama/17736) ... Previously, cmake was forcing `_WIN32_WINNT=0x0A00` for MinGW builds, This caused "macro redefined" warnings with toolchains that define the version. This also removes the `GGML_WIN_VER` variable as it is no longer needed. Signed-off-by: Adrien Gallouët <angt@huggingface.co>	2025-12-12 17:53:16 +02:00
3 people	7adbcafb6c	CUDA: generalized (mma) FA, add Volta support (llama/17505) ... * CUDA: generalized (mma) FA, add Volta support * use struct for MMA FA kernel config --------- Co-authored-by: Aman Gupta <aman>	2025-12-12 17:53:16 +02:00
Tarek DakhranandGeorgi Gerganov	0defeee679	model: LFM2-VL fixes (llama/17577) ... * Adjust to pytorch * Add antialiasing upscale * Increase number of patches to 1024 * Handle default marker insertion for LFM2 * Switch to flag * Reformat * Cuda implementation of antialias kernel * Change placement in ops.cpp * consistent float literals * Pad only for LFM2 * Address PR feedback * Rollback default marker placement changes * Fallback to CPU implementation for antialias implementation of upscale	2025-12-12 17:53:14 +02:00
Georgi Gerganov	968db8bcfa	ggml : add ggml_top_k (llama/17365) ... * ggml : add ggml_top_k * cont : add ggml_argsort_top_k * metal : add top_k support * ggml : cleanup * tests : add virtual err() function for test_case * ggml : add comments	2025-12-12 17:53:08 +02:00
5 people	3e684f26c1	ggml : add ops SOFTPLUS, EXPM1, TRI, SOLVE_TRI, CUMSUM (llama/17063) ... * Add ops needed for new hybrid models: SOFTPLUS, EXPM1, TRI, SOLVE_TRI, CUMSUM * Update ggml/include/ggml.h Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Update tests/test-backend-ops.cpp Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Code review * Whitespace * Update tests/test-backend-ops.cpp Co-authored-by: Diego Devesa <slarengh@gmail.com> * This is actually sigmoid, duh. * Add CONST, remove TRI_KEEP, other changes from review * Update tests/test-backend-ops.cpp Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Update ggml/src/ggml.c Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Update ggml/src/ggml.c Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Update ggml/src/ggml-cuda/unary.cu Co-authored-by: Aman Gupta <amangupta052@gmail.com> * Remove extra script * Update ggml/src/ggml.c Co-authored-by: Diego Devesa <slarengh@gmail.com> * Update tests/test-backend-ops.cpp Co-authored-by: Diego Devesa <slarengh@gmail.com> * moving changes from laptop [no ci] * pre-rebase * Update tests/test-backend-ops.cpp Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com> * Update tests/test-backend-ops.cpp Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com> * Refactor tests * ggml : cleanup * cont : fix ggml_fill srcs * tests : add note * ggml : add ggml_fill_inplace * ggml : add asserts * ggml : fix ggml_fill constant cast * cont : ggml_tri minor * Use TENSOR_LOCALS * Fix regression from #14596, regenerate * Don't make commits at night... --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> Co-authored-by: Diego Devesa <slarengh@gmail.com> Co-authored-by: Aman Gupta <amangupta052@gmail.com> Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>	2025-11-17 21:05:46 +02:00
AclyandGeorgi Gerganov	997fdde0c4	ggml-cpu : bicubic interpolation (llama/16891)	2025-11-09 23:38:03 +02:00
7 people	e1780b209d	model: add support for qwen3vl series (llama/16780) ... * support qwen3vl series. Co-authored-by: Thireus ☠ <Thireus@users.noreply.github.com> Co-authored-by: yairpatch <yairpatch@users.noreply.github.com> Co-authored-by: LETS-BEE <LETS-BEE@users.noreply.github.com> * bugfix: fix the arch check for qwen3vl-moe. * use build_ffn * optimize deepstack structure * optimize deepstack feature saving * Revert "optimize deepstack feature saving" for temporal fix This reverts commit f321b9fdf13e59527408152e73b1071e19a87e71. * code clean * use fused qkv in clip * clean up / rm is_deepstack_layers for simplification * add test model * move test model to "big" section * fix imrope check * remove trailing whitespace * fix rope fail * metal : add imrope support * add imrope support for sycl * vulkan: add imrope w/o check * fix vulkan * webgpu: add imrope w/o check * Update gguf-py/gguf/tensor_mapping.py Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com> * fix tensor mapping --------- Co-authored-by: Thireus ☠ <Thireus@users.noreply.github.com> Co-authored-by: yairpatch <yairpatch@users.noreply.github.com> Co-authored-by: LETS-BEE <LETS-BEE@users.noreply.github.com> Co-authored-by: Xuan Son Nguyen <son@huggingface.co> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>	2025-11-09 23:38:03 +02:00
safranowithandGeorgi Gerganov	757d51d21d	cpu : add FLOOR, CEIL, ROUND and TRUNC unary operators (llama/16083) ... * CPU: Add support for FLOOR,CEIL,ROUND and TRUNC unary operators - Added the operators to unary op enum - Implemented API functions - Implemented forward and unary-op logic in CPU backend - Updated ggml_get_n_tasks - Updated operators names array and static_assert - Updated docs and enabled automatic tests * docs: add documentation for ggml_trunc and ggml_trunc_inplace in ggml.h * chore: remove trailing whitespace from ggml.h * Remove unresolved merge markers * Apply review suggestions: cleanup formatting, enum order and leftover artifacts * Regenerate ops.md using create_ops_docs.py	2025-10-22 12:58:11 +03:00
Reese LevineandGeorgi Gerganov	27ebde6afd	ggml webgpu: add support for soft_max, optimize rms_norm (llama/16357) ... * Add inplace softmax * Move rms_norm to split row approach * Update debug for supports_op * clean up debug statements * Update tests/test-backend-ops.cpp Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>	2025-10-12 11:16:23 +03:00
4 people	33ca8355c4	model : Apertus model implementation (llama/15852) ... * First attempt * No permute during convert (fixes qk tensors), proper norm application. * RoPE = NeoX * Coherence! * Migrate xielu params from tensors to hyperparameters * Simple CUDA kernel * Revert stupid LLM refactorings * Chat template support * configchecker / flake8 errors * Reorder unary.cu * I do conclude that LLMs are, in fact, stupid. * Fix after merge * Final newline * Make xIELU an UNARY_OP * Final newline * Correctly account for parameter shift * Argh. * Update ggml/src/ggml-cpu/unary-ops.cpp Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Refactor: remove unused methods, inline and factorize softplus, add const modifiers * Revert CUDA changes, implement xIELU as a separate OP * Pesky newline * Add float2half / half2float for F16 inputs/outputs * CUDA variants, attempt 2 * Actually, attempt 3 * Update ggml/src/ggml-cuda/unary.cu Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * Missing convert header * Proper formula and reference for xIELU in the comments. * Modify unary-ops.cpp to add the functor-based logic besides the template system to retain optimizations * Apply suggestions from code review Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com> * Add tensor mappings for Apertus to global list instead * Fix lazy on scalars * Update ggml/src/ggml-cuda/unary.cu Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * Add comment about the constraints on positive/negative alpha * Change `softplus` to `ggml_softplus` --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> Co-authored-by: Johannes Gäßler <johannesg@5d6.de> Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>	2025-10-12 11:16:23 +03:00
Reese LevineandGeorgi Gerganov	8208cea829	ggml webgpu: support for rope,div,sub,glu,scale,cont operators (llama/16187) ... * Work on rope * Simplify inplace operation generation and combine mul/add generation * Work on rope variants * implement neox rope * rope complete * Add sub,div,glu operators * implement scale op * Update cpy shader to handle cont/more types * formatting * Update test vars printing for rope,rms_norm * Avoid ROPE hardcoded constants * Add TODO to change ROPE constants to enum Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * fix TODO comment --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>	2025-10-12 11:16:23 +03:00
Georgi Gerganov	eb2c01f92e	metal : refactor + optimize v2 (llama/15995)	2025-09-20 13:46:10 +03:00
Georgi Gerganov	ae6cc6a386	cuda : fix supports_op condition for get_rows when number of blocks is too large (llama/15868) ... * cuda : fix supports_op condition for get_rows when src1->ne2 > 1 ggml-ci * ggml : add comment about ggml_get_rows ggml-ci * cuda : add FIXME [no ci] * cuda : update support condition ggml-ci	2025-09-20 13:42:52 +03:00
Xuan-Son NguyenandGeorgi Gerganov	40bcd1a469	ggml: allow casting between f32 and i32 (llama/15783) ... * ggml: allow casting between f32 and i32 * fix cuda * add vulkan * fix CPU non-cont * add non-cont test case * add note * extend test number range * correct note * add cont version for vulkan	2025-09-20 13:42:51 +03:00
3 people	2228462b19	ggml: add ops for WAN video model (cuda && cpu) (llama/15669) ... * add conv3d support * add ggml_pad_ext for cpu & cuda backend * cuda/cpu: add im2col_3d support * cuda: make im2col a little faster * fix cuda pad/scale/im2col3d * make im2col_3d faster * gguf: support loading tensors which n_dims > GGML_MAX_DIMS * fix cuda get_rows * avoid ggml_conv_3d conflict * correct GGML_OP_COUNT assertion * avoid build failure * avoid build failure on MacOS * cuda: remove unnecessary MIN define * fix cpu im2col_3d * adjust the code style * cuda: use simpler loop in get_rows * add test_im2col_3d to test-backend-ops * test-backend-ops.cpp: remove trailing whitespace * cpu: im2col_3d support non continuous src Co-authored-by: Jeff Bolz <jbolz@nvidia.com> * fix test_im2col_3d * remove unused variables * cuda: get_rows: dfloat2 -> float2 * add test_pad_ext to test-backend-ops.cpp * add gguf_init_from_file_ext impl * Revert "gguf: support loading tensors which n_dims > GGML_MAX_DIMS" This reverts commit d8377a0a37f314bd3713fe043b4333ad661610c1. * Revert "add gguf_init_from_file_ext impl" This reverts commit d9f1d13208c68ef83b3538201ac7f31614fb1994. * update ggml_backend_vk_device_supports_op * fix ggml_backend_vk_device_supports_op * update other backend supports op for ggml_pad_ext * metal/opencl/sycl/vulkan: fix GGML_OP_PAD check in supports_op --------- Co-authored-by: Jeff Bolz <jbolz@nvidia.com>	2025-09-20 13:42:49 +03:00
rmatifandGeorgi Gerganov	d7b7498e76	ggml: add conv3d op (llama/15182) ... * add conv3d * bump GGML_OP_COUNT	2025-09-20 13:42:39 +03:00
R0CKSTARandGeorgi Gerganov	7c077845fd	musa: add GGML_UNUSED_VARS (llama/15446) ... Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com>	2025-09-20 13:42:38 +03:00
4 people	c76ec72d59	finetune: SGD optimizer, more CLI args (llama/13873) ... * examples/finetune -opt SGD (stochastic gradient descent) memory opt add unit tested GGML_OPT_OPTIMIZER_SGD to ggml - avoids allocating m, v tensors. support finetune.cpp arg -opt SGD (or sgd). (default adamw as before) llama 3.2-1b-F32 result: observed 11gb gpu ram (41 sec/epoch) when using SGD instead of 19gb (55 sec/epoch) using adamw. (wikipedia 100 lines finetune) ( using the same GPU memory, adamw can only do before OOM 512 batch/context, reaching: train: [███████▉] data=0000140/0000140 loss=0.02575±0.00099 acc=99.52±0.03% t=00:00:47 ETA=00:00:00 val: [███████▉] data=0000008/0000008 loss=4.76565±0.28810 acc=41.46±0.77% t=00:00:00 ETA=00:00:00 SGD is superior, though it converges slower, with max before OOM 1728 batch/context (esp see the better validation perf): train: [███████▉] data=0000039/0000039 loss=0.00371±0.00010 acc=99.96±0.01% t=00:00:41 ETA=00:00:00 val: [███████▉] data=0000003/0000003 loss=5.11406±0.76034 acc=48.01±0.69% t=00:00:01 ETA=00:00:00 ) note: when finetuning long enough (or w/ enough -lr), validation accuracy *eventually* drops ('catastrophic forgetting') -lr-half (halflife) option useful for SGD to avoid oscillation or super slow underdamped learning (makes setting -lr more forgiving). terminal -lr for now is set by lr-halvings i.e. if you want at most 1/8 the inital -lr you set -lr-halvings 3. note: objective loss not directly comparable between adamw, sgd? - check perplexity or accuracy or consider relative improvements for convergence new finetune args -wd 1e-9 to enable weight decay in sgd or adamw, and max -epochs N (default 2 as before) cache (1 - wd*alpha) in 'adamw' opt struct - no noticeable perf benefit, disabled (still done for new SGD though) since opt. memory is pre-allocated, the ggml_opt_get_optimizer_params would probably be able to change between SGD and AdamW with each epoch but would need to use adamw for the first (unconfirmed - no cmdline arg to set such a policy yet) test-opt checks adamw as before and now sgd (except for a few disabled tests for sgd only; probably just needs logging values and adding alternate reference values); tolerance on the 'regression' test is broader for sgd (so we don't need many more epochs) * Vulkan: Implement GGML_OP_OPT_STEP_SGD * tests: Fix OPT_STEP_SGD test-backend-ops * SGD op param store weight-decay and not 1-alpha*wd * minor + cosmetic changes * fix vulkan sgd * try CI fix --------- Co-authored-by: 0cc4m <picard12@live.de> Co-authored-by: Johannes Gäßler <johannesg@5d6.de>	2025-08-18 20:30:45 +03:00
JuddandGeorgi Gerganov	80ef57f0f0	ggml : update ggml_rope_multi (llama/12665) ... * update `rope_multi`: 1. add `ggml_rope_multi_inplace`; 1. use `GGML_MROPE_SECTIONS` instead of 4. * Apply suggestions from code review Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>	2025-08-18 20:30:45 +03:00
3 people	d3aab3efde	llama : add gpt-oss (llama/15091) ... * oai moe * compat with new checkpoint * add attn sink impl * add rope scaling yarn * logits match with latest transformers code * wip chat template * rm trailing space * use ggml_scale_bias * rm redundant is_swa_all * convert interleaved gate_up * graph : fix activation function to match reference (llama/7) * vocab : handle o200k_harmony special tokens * ggml : add attention sinks support (llama/1) * llama : add attn sinks * ggml : add attn sinks * cuda : add attn sinks * vulkan : add support for sinks in softmax remove unnecessary return * ggml : add fused swiglu_oai op (llama/11) * ggml : add fused swiglu_oai op * Update ggml/src/ggml-cpu/ops.cpp Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * update CUDA impl * cont : metal impl * add vulkan impl * test-backend-ops : more test cases, clean up * llama : remove unfused impl * remove extra lines --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> --------- Co-authored-by: slaren <slarengh@gmail.com> * repack mxfp4 upon conversion * clean up a bit * enable thinking * add quick hack to render only some special tokens * fix bf16 conversion * remove vocab hack * webui ok * support chat parsing for gpt-oss * fix webui * direct mapping mxfp4, FINALLY * force using mxfp4 * properly use lazy tensor * ggml : add mxfp4 ggml : use e8m0 conversion instead of powf Co-authored-by: Diego Devesa <slarengh@gmail.com> change kvalues_mxfp4 table to match e2m1 (llama/6) metal : remove quantization for now (not used) cuda : fix disabled CUDA graphs due to ffn moe bias vulkan : add support for mxfp4 cont : add cm2 dequant * ggml : add ggml_add_id (llama/13) * ggml : add ggml_add_id * add cuda impl * llama : add weight support check for add_id * perf opt * add vulkan impl * rename cuda files * add metal impl * allow in-place ggml_add_id * llama : keep biases on CPU with --cpu-moe * llama : fix compile error ggml-ci * cuda : add fallback for __nv_cvt_e8m0_to_bf16raw ggml-ci * cleanup ggml-ci * sycl : fix supports_op for MXFP4 ggml-ci * fix Unknown reasoning format * ggml-cpu : fix AVX build ggml-ci * fix hip build ggml-ci * cuda : add mxfp4 dequantization support for cuBLAS ggml-ci * ggml-cpu : fix mxfp4 fallback definitions for some architectures ggml-ci * cuda : fix version required for __nv_cvt_e8m0_to_bf16raw --------- Co-authored-by: Xuan Son Nguyen <son@huggingface.co> Co-authored-by: slaren <slarengh@gmail.com>	2025-08-18 20:30:45 +03:00
Georgi Gerganov	85dcc74b88	sync : resolve conflicts (ggml/0) ... ggml-ci	2025-07-12 19:23:56 +03:00
Xuan-Son NguyenandGeorgi Gerganov	2021870fb8	ggml : add ggml_scale_bias (llama/14417) ... * ggml : add ggml_scale_bias * ggml_vec_mad1_f32 * add more simd * add CUDA * sycl * vulkan * cann (placeholder) * opencl * will this fix cpu? * fix cuda * suggestions from coderabbit * fix cann compile error * vDSP_vsmsa * rm __ARM_FEATURE_SVE * use memcpy for op params * make code looks more consistent * use scalar for __ARM_FEATURE_SVE * add x param to ggml_vec_mad1_f32	2025-07-12 19:23:56 +03:00
Aman GuptaandGeorgi Gerganov	98e7beac6c	CUDA: add bilinear interpolation for upscale (llama/14563)	2025-07-12 19:23:56 +03:00
Sigbjørn SkjæretandGeorgi Gerganov	e8138c51d2	ggml : implement GEGLU_ERF and GEGLU_QUICK ops (llama/14445)	2025-07-12 19:23:56 +03:00
Georgi Gerganov	0abd0660e1	ggml : fix FA mask dim 2 and 3 (llama/14505) ... * ggml : fix FA mask dim 2 and 3 ggml-ci * backends : unsupport batched FA in CUDA and Vulkan ggml-ci * vulkan : disable FA for mask->ne[2] != 1	2025-07-12 19:23:56 +03:00
compiladeandGeorgi Gerganov	d2d120c256	llama : initial Mamba-2 support (llama/9126) ... * llama : initial Mamba-2 support * ggml : SIMD ggml_ssm_scan for Mamba-2 * ggml : improve ggml_mul speed when masking recurrent states * llama : support running Mamba-Codestral-7B-v0.1 * llama : fix Mamba-2 conv state saving * ggml : make the ggml_mul fast broadcast path more consistently formatted * llama : remove unused variable * llama : add missing break * convert_hf : prefer SentencePiece tokenizer for Mamba-2 when present The tokenzier.json of Mamba-Codestral-7B-v0.1 otherwise requires workarounds to work correctly. * llama : avoid redundant state copy for Mamba 1 and 2 * metal : attempt to adapt SSM_SCAN for Mamba-2 * metal : fix SSM_SCAN pipeline scope * metal : use log and exp instead of log1pf and expf in SSM_SCAN * metal : remove unused arguments for SSM_SCAN The max index is 31, so trimming the arguments is necessary. * metal : add back n_seqs to SSM_SCAN args Whoops, this is needed for the offset in the concatenated output. * metal : fix SSM_SCAN state head offset * metal : fix wrong number of tokens per sequence in SSM_SCAN * ggml : remove unused fast broadcast path in GGML_MUL This was initially added because states were masked with ggml_mul, but this is no longer done and so this "optimisation" is no longer necessary, or at least not worth the additional code complexity. * ggml : avoid multiply by D in GGML_OP_SSM_SCAN This makes the weight buft detection in src/llama.cpp simpler. * convert : transpose Mamba-2 A, D and reshape SSM_NORM This breaks existing conversions of Mamba-2 models to avoid some reshapes. Not sure if it's a good idea, but it makes the graph slightly cleaner. * llama : more appropriate SSM_SCAN and SSM_CONV buft support checks * convert : fix flake8 lint * metal : fix confusion between ; and , * metal : add missing args for nb references in ssm_scan_f32_group * metal : single-user mamba2 inference works * kv-cache : remove const_cast when setting inputs for s_copy And also fix multi-user inference for recurrent models by using cell_id instead of i as the kv cell index when populating s_copy. * convert : avoid AutoConfig for Mamba and Mamba2 hparams * kv-cache : allow context shift for recurrent models * graph : fix recurrent state copies when avoiding copies Works, but using lambda functions might not be that clean. * ggml : fix mamba2 ssm scan when compiled with SVE * ggml-cpu : reorder SVE FMA for consistency with other SIMD arches * cuda : implement ssm scan for Mamba2 There is still room for improvement, but it works! * cuda : adapt Mamba1 ssm scan to shape changes from Mamba2 * mamba : fix mismatched new and delete size for llm_build_mamba Subclasses of llm_graph_context cannot have extra fields, because the called destructor is not the one from the subclass. This otherwise would cause problems when runnning Mamba-(1|2) inference when compiled -DGGML_SANITIZE_ADDRESS=ON * cuda : graceful fallback for Mamba-1 models with weird embd size	2025-07-12 19:23:56 +03:00
Georgi Gerganov	d1286cf32b	ggml : support bcast ggml_soft_max_ext, ggml_flash_attn_ext (llama/14435)	2025-07-12 19:23:56 +03:00