* 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>
* hex-mm: new weight layout and fusion updates
* hvx-mm: unroll the new tiled vec_dots to optimize hvx register util
* hex-mm: optimize dyn.quant format for q8_0 and q8_1 to reduce overhead in vec_dots.
* hvx-mm: parallel quantizer per block for large rows
* hvx-mm: simplify and futher optimize dyn.quant and vec_dots
* hvx-mm: keep intermediate per tile accumulators in fp16
* hmx-mm: optimize weight dequant by aligning the repacked tiles with the DMA
* hmx-mm: remove qweight scratch and just use vtcm_weight
* hmx-mm: remove all unused and obsolete code
* hmx-mm: the new tiled repack format is here to stay -- rename all x4x2 to _tiled
* hmx-mm: improve activation processing with dma prefetch
* hex-mm: fix hmx/hvx fallback logic and MUL_MAT_ID allocation (unbreaks OLMoE)
* hex-mm: align the weight tiles with dma just like we did in hmx-mm
* hex-mm: factor out common mm bits into htp/matmul-ops.h
* hex-mm: start moving mm kernel selection to the host
* hex-mm: move all of the matmul param compute into the host
* hmx-mm: restore pipelined mode
* hmx-mm: unroll the dequant functions to optimize register usage
* hmx-mm: further improve activation process
* hex-mm: use vtcm_seq_alloc for all vtcm allocations and define more common functions
* hex-mm: improve mm optimizer to acount for number of activation threads
* hex-mm: fix matmul-id kernel params selection (unbreaks OLMoE and LFM)
* hexagon: remove support for arch < v73 since HMX is now required for most use-cases
* hex-mm: cleanup naming for consistency
* hex-mm: make sure matmul fusion accounts for vtcm allocation
* hex-mm: minor cleanup for kernel_params definition
* hex-mm: replace hardcoded limits with proper checks for vtcm requirements
* hex-mm: add support for non-tiled mm as a fallback option and factor out hvx kernels into separate header
* hex-mm: remove unused functions
* hex-mm: add shorthand for MM_SELECT in run-tool script
* hvx-mm: factor out hvx/hmx microkernels and unify matmul entry and dispatch
* hex-mm: further cleanup matmul fallback path
* hex-mm: refactor matmul entry point and dispatch a bit further
* hexagon: update cmake build to enable hmx for everything
* hex-ops: optimize kernel_param updates and include summary in the logs
* hex-mm: add support for GGML_HEXAGON_MM_SELECT
* hex-mm: add hex-common header
* hex-mm: pass correct number of tasks to workpool
* hex-mm: add proper checks for no-work in dyn.quant tasks
* hex-mm: convert all quantizers into a macro
* hex-mm: fix hvx-flat fallback to pass all MUL_MAT tests
* hex-mm: vectorize q8_1 quantizer
* hex-mm: improve fused ffn mm stride handling
* hex-mm: consistent use of n_threads and pipeline in kernel_params
* hexagon: minor formatting
* hex-mm: update MUL_MAT_ID kernel_param handling to make sure host/npu are in sync
* hvx-mm: go back to accumulating in fp32 in tiled hvx kernels, more accurate and same perf
* hvx-mm: unroll the loops and remove masking that is not needed for tiled accums
* hmx-mm: optimize activation processing (slit loops, some unrolling, etc)
* hmx-mm: minor optimization for output processing
* hex-mm: consistent use of uint32_t and size_t in mm kernels
* hex-mm: remove legacy restrictions for rows to be multiple of 256
* hexagon: replace sprintf with snprintf
* hex-mm: relax hardcoded nrows checks and rely on VTCM size requirements
* hexagon: minor alignment fix
* hexagon: fix trailing spaces
* hex-mm: relax padding from 256 to 128 (leftovers)
* hex-mm: remove redundant checks for weight align to 128
we always use 2D dma for the weights and align them properly
* hmx-mm: MUL_MAT_ID better work distribution between hvx threads and hmx tracing
* hex-mm: specialize per-token mmid activation handling
* hex-profile: update python scripts to handle kernel-params section in the logging output
* hex-mm: move n_prefetch (aka dma_depth) into kernel params and remove unused fields
* hex-trace: use easier to parse format, simply and fix post-proc scripts
* hmx-mm: relax 32 row limit for output processing which helps utilization
* hmx-mm: use start-chunk idx for tracing info
* hmx-mm: parameterize activation dma pipeline
* hexagon: add support for simple graph caching to avoid recomputing kernel-params
* hex-mm: remove left-over repack functions
* hex-mm: tighten n_prefetch asserts
* hex-mm: remove duplicate round/align_up helper
* hexagon: cleanup common header used in host/npu
* hexagon: update early wakeup threshold
* hmx-mm: define cost constants and update solver to assume that repacked ne[1] is padded to 32
* hmx-mm: make precompute_matmul a bit more readable (split into smaller functions, etc)
* hex-mm: remove n_threads constraint
* hex-mm: minor formatting updates
* hex-mm: remove obsolete profiling logs
* hex-mm: restore hardcode gate to refuse lm-head to avoid repacking that tensor
* rename GGML_SYCL_SUPPORT_LEVEL_ZERO to GGML_SYCL_SUPPORT_LEVEL_ZERO_API, and GGML_SYCL_ENABLE_LEVEL_ZERO to GGML_SYCL_USE_LEVEL_ZERO_API
* fix code format
* fix error when rebase
* SYCL: fix multi-GPU system RAM exhaustion by using Level Zero allocations
Replace sycl::malloc_device with zeMemAllocDevice for GPU memory allocation
in the SYCL backend. sycl::malloc_device triggers the xe kernel driver's
DMA-buf/TTM path which mirrors every VRAM allocation 1:1 in system RAM.
zeMemAllocDevice uses the SVM/P2P path with no host staging.
On a dual Intel Arc Pro B70 system (64GB VRAM, 64GB RAM), a 15.6 GiB model
consumed 60 GiB of system RAM via sycl::malloc_device, causing OOM crashes.
With zeMemAllocDevice, the same workload uses ~6.7 GiB of system RAM with
no performance regression.
All Level Zero calls include automatic fallback to the original SYCL
allocation path if Level Zero interop is unavailable.
* SYCL: address review feedback - remove try/catch, check device types, deduplicate
- Remove try/catch from malloc/free/memcpy helpers, check backend and
device type upfront instead (ggml_sycl_is_level_zero, ggml_sycl_is_dgpu)
- Move shared helpers (is_level_zero, is_dgpu, free_device) to common.cpp
and declare in common.hpp to eliminate code duplication
- Use SYCL_CHECK(CHECK_TRY_ERROR()) for fallback sycl::free calls
- Guard dev2dev_memcpy L0 path to dGPU-to-dGPU only, preserving the
host-staged path for iGPU-to-dGPU transfers
- Add Windows Level Zero SDK path detection (LEVEL_ZERO_V1_SDK_PATH)
in CMakeLists.txt (co-authored with @arthw)
* SYCL: add build/runtime flags for Level Zero, address review feedback
Implements the architecture suggested by @arthw: compile-time and runtime
flags to cleanly separate Level Zero and SYCL memory API paths.
- Add GGML_SYCL_SUPPORT_LEVEL_ZERO cmake option (default ON). All Level
Zero code is wrapped in #ifdef so the build works on systems without
the Level Zero SDK installed (e.g. CPU-only CI servers). Both the
loader library and headers are checked before enabling.
- Add GGML_SYCL_ENABLE_LEVEL_ZERO runtime env var (default 1). Controls
whether Level Zero or SYCL memory APIs are used. Only one API style is
used per session, no mixing. If Level Zero is enabled but the devices
don't support the Level Zero backend, it auto-disables with a warning.
- Remove Level Zero code from dpct_malloc. It was unused (dpct::device_memory
is not called anywhere in the backend) and used try/catch for flow control.
- Update SYCL.md with documentation for both new parameters.
Tested on Intel Arc Pro B70 (32GB), single-GPU and dual-GPU, with both
GGML_SYCL_SUPPORT_LEVEL_ZERO=ON and OFF builds. AI-assisted development
(Claude). Code reviewed and tested on my hardware.
* SYCL: unify Level Zero malloc/free call sites, address review feedback
Move ggml_sycl_malloc_device to common.cpp alongside ggml_sycl_free_device.
Both functions are now unconditionally available — Level Zero code is
#ifdef'd inside the functions, not at call sites. All call sites use
uniform SYCL_CHECK(CHECK_TRY_ERROR()) wrapping with no #ifdef blocks.
Addresses arthw's review: wrap all malloc/free in SYCL_CHECK for stack
traces on failure, eliminate duplicated #ifdef/else patterns at 6 call
sites (-29 lines net).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* SYCL: add Level Zero SDK to CI, fix device check and missed alloc paths
Add Level Zero SDK installation to Ubuntu and Windows SYCL CI jobs
so the Level Zero code path is compiled and tested in CI.
Fix two bugs found during extended dual-GPU testing (no
ONEAPI_DEVICE_SELECTOR set):
- The Level Zero backend check was iterating all SYCL devices
including CPU. The OpenCL CPU device caused Level Zero to be
disabled for the GPUs, defeating the fix on multi-GPU systems.
Added is_gpu() filter so only GPU devices are checked.
- sycl_ext_malloc_device/sycl_ext_free (tensor reorder temp buffers)
were still calling sycl::malloc/sycl::free directly, bypassing the
Level Zero path. Routed through ggml_sycl_malloc_device/free_device
for consistency with the other device memory call sites.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* SYCL: address arthw review feedback on Level Zero memory API structure
- Move ggml_sycl_malloc_device to static function in ggml-sycl.cpp;
only ggml_sycl_free_device (used by common.cpp) stays in common.cpp
- Switch both helpers to use g_ggml_sycl_enable_level_zero global
instead of per-call queue backend checks
- Remove #ifdef wrapper from global definition; always declare at 0,
add #else branch in init block so it stays 0 when L0 not compiled in
- Update init loop comment to explain GPU-only device check
- CMakeLists: message(STATUS) before the if block; align option wording
AI-assisted implementation. Reviewed and tested on dual Intel Arc Pro
B70 (32 GB each): test-backend-ops OK on both GPUs, single/dual-GPU
Q4_K_M and Q8_0 bench correct, zeMemAllocDevice GTT delta confirmed
<5 MiB per 4 GiB allocation (vs ~4 GiB shadow with sycl::malloc_device).
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
* SYCL: remove unused cstdio/cstdlib includes from common.cpp
Leftover from the deleted ggml_sycl_queue_supports_level_zero helper.
Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
* Apply suggestions from code review
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* SYCL: preserve Level Zero allocation path during early malloc
* ci: fix Level Zero package conflict in Intel Docker build
* ci: find Level Zero loader in oneAPI package step
* ci: allow Windows SYCL package without Level Zero DLL
---------
Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
In #11362 hip graph was disabled by default as, at the time, its performance impact was negative. Due to improvements in rocm and our usage and construction of graphs this is no longer true, so lets change the default.
#21630 added the CMP0194 NEW policy to silence a CMake warning, but on Windows runners it caused CMake to prefer the MinGW toolchain for ASM and broke MSVC builds.
Reverting only that policy block restores the previous working behavior. The CMake 4.1+ warning comes back, but that is cosmetic and does not break any platform.
Reported-by: oobabooga
Refs: #21630
Co-authored-by: texasich <texasich@users.noreply.github.com>
* [SYCL] Fix Q8_0 reorder: add missing dequantize path for GEMM
The Q8_0 reorder optimization (#21527) was missing a reorder-aware
dequantizer for the GEMM code path used during prompt processing.
After token generation reordered Q8_0 weights (via DMMV/MMVQ), the
next prompt processing pass would read them with the standard
dequantizer, producing garbage output.
Add dequantize_block_q8_0_reorder() and wire it into both
ggml_get_to_fp16_sycl() and ggml_get_to_fp32_sycl(), matching the
pattern already used by Q4_0, Q4_K, and Q6_K.
Fixes#21589
AI (Claude) was used to assist with root cause investigation and
writing the kernel code. All code was human-reviewed and tested
on real hardware.
* SYCL: fix reorder crash when device memory is full
The reorder optimization allocates a temporary buffer the full size of
the weight tensor on the device. When VRAM is nearly full (large models
on a single GPU), this allocation fails and the subsequent memcpy crashes
on a NULL pointer.
Fix: try device allocation first, fall back to host memory if device
memory is full. The reorder kernel still works correctly reading from
host memory over PCIe. This is slower for the one-time reorder (~21 t/s
vs ~38 t/s on Intel Arc Pro B70), but the optimization is preserved for
all subsequent inference. If both device and host allocation fail, skip
the reorder and fall back to the unoptimized kernel path.
Also fixes a bug where opt_for_reorder() marked tensors as reordered
even when the reorder was skipped due to allocation failure. This caused
DMMV/MMVQ kernels to read the original AoS data as if it were SoA,
producing garbage output or NaN results.
Tested on Intel Arc Pro B70 (32GB) with Q8_0, Q4_K_M models. Coding was
AI-assisted (Claude), reviewed and tested on hardware by a human.
Fixes#20478
* SYCL: add RAII temp buffer class + macro guard for host fallback
Replace sycl_ext_malloc_with_fallback/sycl_ext_free_fallback free
functions with sycl_reorder_temp_buffer RAII class. The host_fallback
bool is now a private member, and cleanup happens automatically at
scope exit.
Add GGML_SYCL_HOST_MEM_FALLBACK cmake option (default ON) to guard
the host memory fallback code path. Device access to host memory
requires Linux kernel 6.8+ (Ubuntu 26.04+); users on older kernels
can set -DGGML_SYCL_HOST_MEM_FALLBACK=OFF to disable it.
Addresses arthw's review on PR #21638.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* SYCL: document GGML_SYCL_HOST_MEM_FALLBACK build option in SYCL.md
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* SYCL: add reorder-aware DMMV dequantizers for Q4_K and Q6_K
Q4_K and Q6_K had reorder support for MMVQ and GEMM paths but not
DMMV. When the DMMV path encountered reordered data it would abort.
Add DMMV kernels that read from the SOA reorder layout for both
types. Same math as the non-reorder versions, different memory
access pattern.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* cmake: fix CMP0194 warning on Windows with MSVC
Set CMP0194 policy to NEW before project() call in ggml/CMakeLists.txt to suppress the "MSVC is not an assembler for language ASM" warning introduced in CMake 4.1.
The ggml project enables ASM globally for Metal (macOS) and KleidiAI (ARM) backends. On Windows/MSVC, no assembler sources are used, but CMake 4.1+ warns because cl.exe is not a valid ASM compiler.
This follows the same pattern used in ggml-vulkan (CMP0114, CMP0147).
Closesggml-org/llama.cpp#20311
* cmake: apply cisc's formatting suggestion
---------
Co-authored-by: texasich <texasich@users.noreply.github.com>
* ggml: backend-agnostic tensor parallelism
* support for GPT-OSS, Qwen 3 MoE
* partial Vulkan fix
* add support for 4/8 GPUs
* unconditional peer access
* re-use buffers + ggml contexts
* fix output pattern
* NCCL support
* GGML: HIP: add RCCL support
* Remove shfl and AllReduce from backend interface
* move allocation workaround out of ggml-alloc.c
* 2d tensor set/get support
* Fix the seg fault without NCCL
* Apply suggestion from JohannesGaessler
* support for tensor dims % n_devs != 0
* fix view_offs scaling
* arbitrary num. of GPUs/tensor split
* fix compilation
* better granularity estimate
* Support device-specific host buffer types if all underlying backends expose the same type. This allows using pinned memory instead of pageable memory for CUDA.
Fix compilation errors.
* partial Qwen 3 Next support
* Fix qwen3 30b (llama/8)
* Fix crash with Qwen-30B-A3B Q4_0
Qwen-30B-A3B Q4_0 has an intermediate dimension of 768. Using a granularity of 256 forces an uneven split between GPUs, which is not supported by the current implementation.
* Decide block size based on tensor quantization type
* Fix crashes due to KV cache serialization (llama/9)
KV cache serialization requires non-zero offsets on the tensor. Add support in the meta backend to set/get a tensor with a non-zero offset.
* metal : fix build (llama/7)
* static memory allocations, fix usage count
* fix tensor granularity
* more even memory distribution
* use BF16 for allreduce
* rebase fixup
* better error message for unsupported architectures
* Fix device mismatch during scatter of allReduce. (llama/11)
There is a mismatch between the dst buffer device and the backend device, causing the use of sync copies
* Enable the previous allreduce implementation. It is better in both perf and stability (llama/12)
* delay AllReduce for Moe for less I/O
* build : clean-up compile warnings
* backend : move most of the meta backend API to ggml-backend-impl.h
* cont : hide unused public API in the implementation
* llama : use llama_device + remove ggml_backend_dev_is_meta()
* ggml-backend : remove unused alloc include
* minor : remove regex include
* ggml : introduce ggml-ext.h for staging new APIs
* rebase fixup
* fix tests
* llama : more robust logic for determining Meta devices (llama/16)
* llama : more robust logic for determining Meta devices
* cont : fix devs size check
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* cont : fix log type
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* disable roundtrip for meta backend
* fix arch selection
* Qwen 3.5 support
* fix Gemma 4 MoE
* fix OpenVino, SYCL
* fix test-llama-archs for CPU-only builds
* Fix Qwen 3.5 MoE
* disable meta backend tests for WebGPU
* tests : filter CPU-based devices from the Meta backend tests (llama/17)
* meta : formatting, naming, indentation (llama/18)
* formatting : llama-model.cpp
* formatting : ggml-ext.h
* formatting : ggml-backend-meta.cpp
* meta : add TODO
* add documentation
* better error messages
* fix GPT-OSS
---------
Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz>
Co-authored-by: Gaurav Garg <gaugarg@nvidia.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* Update build doc
* Add cgraph tensor output name to OV op name
* Update openvino build instructions
* Add initial NPU support
* draft NPU support version 2: prefill + kvcache
* NPU support version 2: prefill + kvcache
* Change due to ggml cgraph changes, not correct yet
* Change due to ggml cgraph changes, llama-3.2 CPU work
* Add AMD64 to CMakeLists
* Change due to ggml cgraph changes, all device work
* Refactor: clean, fix warning
* Update clang-format
* Statful transformation for CPU GPU
* Add SwiGLU
* Fuse to SDPA
* Replace Concat with Broadcast in MulMat for GQA
* Pull out indices creation for kv cache update
* Refactor: remove past_token_len from extra_inputs
* Fix Phi3 SwiGLU and SoftMax
* Pull out sin cos from rope
* Reduce memory: free ov weights node after graph conversion
* Fix CPY due to cgraph change
* Added OpenVINO CI/CD. Updated docs
* Fix llama-cli
* Fix Phi3 ROPE; Add test-backend-ops
* Fix NPU
* Fix llama-bench; Clang-format
* Fix llama-perplexity
* temp. changes for mark decomp
* matmul in fp32
* mulmat input conversion fix
* mulmat type conversion update
* add mark decomp pass
* Revert changes in fuse_to_sdpa
* Update build.md
* Fix test-backend-ops
* Skip test-thread-safety; Run ctest only in ci/run.sh
* Use CiD for NPU
* Optimize tensor conversion, improve TTFT
* Support op SET_ROWS
* Fix NPU
* Remove CPY
* Fix test-backend-ops
* Minor updates for raising PR
* Perf: RMS fused to OV internal RMS op
* Fix after rebasing
- Layout of cache k and cache v are unified: [seq, n_head, head_size]
- Add CPY and FLASH_ATTN_EXT, flash attn is not used yet
- Skip test-backend-ops due to flash attn test crash
- Add mutex around graph conversion to avoid test-thread-safety fali in the future
- Update NPU config
- Update GPU config to disable SDPA opt to make phi-3 run
* Change openvino device_type to GPU; Enable flash_attn
* Update supports_buft and supports_op for quantized models
* Add quant weight conversion functions from genai gguf reader
* Quant models run with accuracy issue
* Fix accuracy: disable cpu_repack
* Fix CI; Disable test-backend-ops
* Fix Q4_1
* Fix test-backend-ops: Treat quantized tensors as weights
* Add NPU Q4_0 support
* NPU perf: eliminate zp
* Dequantize q4_1 q4_k q6_k for NPU
* Add custom quant type: q8_1_c, q4_0_128
* Set m_is_static=false as default in decoder
* Simpilfy translation of get_rows
* Fix after rebasing
* Improve debug util; Eliminate nop ReshapeReshape
* STYLE: make get_types_to_requant a function
* Support BF16 model
* Fix NPU compile
* WA for npu 1st token acc issue
* Apply EliminateZP only for npu
* Add GeGLU
* Fix Hunyuan
* Support iSWA
* Fix NPU accuracy
* Fix ROPE accuracy when freq_scale != 1
* Minor: not add attention_size_swa for non-swa model
* Minor refactor
* Add Q5_K to support phi-3-q4_k_m
* Requantize Q6_K (gs16) to gs32 on GPU
* Fix after rebasing
* Always apply Eliminate_ZP to fix GPU compile issue on some platforms
* kvcachefusion support
* env variable GGML_OPENVINO_DISABLE_SDPA_OPTIMIZATION added
* Fix for Phi3
* Fix llama-cli (need to run with --no-warmup)
* Fix add_sliced_mask; Revert mulmat, softmax; Remove input attention_size, iSWA model not working
* fix after rebasing
* Fix llama-3-8b and phi3-mini q4_0 NPU
* Update to OV-2025.3 and CMakeLists.txt
* Add OV CI cache
* Apply CISC review and update CI to OV2025.3
* Update CI to run OV dep install before build
* Update OV dockerfile to use OV2025.3 and update build docs
* Style: use switch in supports_ops
* Style: middle ptr and ref align, omit optional struct keyword
* NPU Unify PD (llama/14)
* Stateless. Fix llama-cli llama-server
* Simplify broadcast op in attention
* Replace get_output_tensor+memcpy with set_output_tensor
* NPU unify PD. Unify dynamic and static dims
* Clean placeholders in ggml-openvino.cpp
* NPU unify PD (handled internally)
* change graph to 4d, support multi sequences
* Fix llama-bench
* Fix NPU
* Update ggml-decoder.cpp
Hitting error while compiling on windows:
error C3861: 'unsetenv': identifier not found
Reason: unsetenv() is a POSIX function; it doesn’t exist on Windows. Visual Studio (MSVC) won’t recognize it.
Proposed fix: Use _putenv_s() (Windows equivalent)
This is supported by MSVC and achieves the same effect: it removes the environment variable from the process environment.
This keeps cross-platform compatibility.
* Update ggml-decoder.cpp
* Update ggml-decoder.cpp
* Update ggml-decoder.cpp
* Update ggml-decoder.cpp
* Update ggml-decoder.cpp
* Remove the second decoder for node. Moving the function into the model decoder
* Fix error for naive
* NPU prefill chunking
* NPU fix llama-bench
* fallback naive run with accuracy issue
* NPU support llma-perplexity -b 512 --no-warmup
* Refactor: split ov_graph_compute for dynamic and static
* remove unused API GgmlOvDecoder::get_output_stride(const std::string & name)
* minor update due to ov 2025.4
* remove unused API GgmlOvDecoder::get_output_names()
* remove unused API get_output_shape(const std::string & name)
* Modified API GgmlOvDecoder::get_output_type(const std::string & name)
* Removed API GgmlOvDecoder::get_output_op_params(const std::string & name)
* Removed API get_output_ggml_tensor(const std::string & name)
* Removed API m_outputs
* Removed m_output_names
* Removed API GgmlOvDecoder::get_input_names()
* Removed API GgmlOvDecoder::get_input_stride(const std::string& name)
* Removed API get_input_type
* Removed API get_input_type
* Removed API GgmlOvDecoder::get_input_shape(const std::string & name)
* Removed API GgmlOvDecoder::get_input_op_params(const std::string & name)
* Fix error for decoder cache
* Reuse cached decoder
* GPU remove Q6_K requantization
* NPU fix wrong model output shape
* NPU fix q4 perf regression
* Remove unused variable nodes
* Fix decoder can_reuse for llama-bench
* Update build.md for Windows
* backend buffer: allocate on host
* Use shared_buffer for GPU NPU; Refactor
* Add ov_backend_host_buffer; Use cached remote context
* Put kvcache on GPU
* Use ggml_aligned_malloc
* only use remote tensor for kvcache
* only use remote tensor for kvcache for GPU
* FIX: use remote tensor from singleton
* Update build.md to include OpenCL
* NPU always requant to q4_0_128
* Optimize symmetric quant weight extraction: use single zp
* Use Q8_0_C in token embd, lm_head, and for 5 and 6 bits quant
* Update build.md
* Support -ctk f32
* Initial stateful graph support
* Update ggml/src/ggml-openvino/ggml-decoder.cpp
Co-authored-by: Yamini Nimmagadda <yamini.nimmagadda@intel.com>
* code cleanup
* npu perf fix
* requant to f16 for Q6 embed on NPU
* Update ggml/src/ggml-openvino/ggml-decoder.cpp
* Update ggml/src/ggml-openvino/ggml-openvino-extra.cpp
* Create OPENVINO.md in llama.cpp backend docs
* Update OPENVINO.md
* Update OPENVINO.md
* Update OPENVINO.md
* Update build.md
* Update OPENVINO.md
* Update OPENVINO.md
* Update OPENVINO.md
* kq_mask naming fix
* Syntax correction for workflows build file
* Change ov backend buffer is_host to false
* Fix llama-bench -p -n where p<=256
* Fix --direct-io 0
* Don't put kvcache on GPU in stateful mode
* Remove hardcode names
* Fix stateful shapes
* Simplification for stateful and update output shape processing
* Remove hardcode names
* Avoid re-compilation in llama-bench
* Extract zp directly instead of bias
* Refactor weight tensor processing
* create_weight_node accept non-ov backend buffer
* remove changes in llama-graph.cpp
* stateful masking fix (llama/38)
Fix for stateful accuracy issues and cl_out_of_resources error in stateful GPU with larger context sizes.
* Fix test-backend-ops crash glu, get_rows, scale, rms_norm, add
* hardcoded name handling for rope_freqs.weight
* Suppress logging and add error handling to allow test-backend-ops to complete
* Fix MUL_MAT with broadcast; Add unsupported MUL_MAT FLASH_ATTN cases
* Use bias instead of zp in test-backend-ops
* Update OV in CI, Add OV CI Tests in GH Actions
* Temp fix for multithreading bug
* Update OV CI, fix review suggestions.
* fix editorconfig-checker, update docs
* Fix tabs to spaces for editorconfig-checker
* fix editorconfig-checker
* Update docs
* updated model link to be GGUF model links
* Remove GGML_CPU_REPACK=OFF
* Skip permuted ADD and MUL
* Removed static variables from utils.cpp
* Removed initializing non-existing variable
* Remove unused structs
* Fix test-backend-ops for OV GPU
* unify api calling
* Update utils.cpp
* When the dim is dynamic, throw an error, need to is stastic forst
* Add interface compute_model_outputs(), which get the model output through computing the node use count & status in the cgraph to avoid the flag using
* No need to return
* Fix test-backend-ops for OV GPU LNL
* Fix test-thread-safety
* use the shape from infer request of output tensor create to avoid issue
* fix dynamic output shape issue
* fix issue for the unused node in tests
* Remove unused lock
* Add comment
* Update openvino docs
* update to OV release version 2026.0
* add ci ov-gpu self hosted runner
* fix editorconfig
* Fix perplexity
* Rewrite the model inputs finding mechanism (llama/54)
* Rewrite the model inputs finding logistic
* Put stateful shape handle in get input shape
* Put the iteration logistic in func
* Added ggml-ci-intel-openvino-gpu and doc update
* .hpp files converted to .h
* fix ggml-ci-x64-intel-openvino-gpu
* Fix for stateful execution bug in llama-bench
* Minor updates after stateful llama-bench fix
* Update ggml/src/ggml-openvino/utils.cpp
Co-authored-by: Yamini Nimmagadda <yamini.nimmagadda@intel.com>
* Remove multiple get_shape calls
* Bring back mutex into compute
* Fix VIEW op, which slice the input node
* Added token_len_per_seq existence check before slicing masks and moved node retrieval inside guarded block to prevent missing-key access
* Temp. fix for test requant errors
* Update to OV ggml-ci to low-perf
* ci : temporary disable "test-llama-archs"
* ci : cache v4 -> v5, checkout v4 -> v6, fix runner tag
* docs : update url
* Fix OV link in docker and Update docs
---------
Co-authored-by: Ravi Panchumarthy <ravi.panchumarthy@intel.com>
Co-authored-by: Cavus Mustafa <mustafa.cavus@intel.com>
Co-authored-by: Arshath <arshath.ramzan@intel.com>
Co-authored-by: XuejunZhai <Xuejun.Zhai@intel.com>
Co-authored-by: Yamini Nimmagadda <yamini.nimmagadda@intel.com>
Co-authored-by: Xuejun Zhai <Xuejun.Zhai@intel>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Some backend depends on CMAKE_RUNTIME_OUTPUT_DIRECTORY to create temporary file like metal backened.
Missing CMAKE_RUNTIME_OUTPUT_DIRECTORY will cause some cmake error like permission denied (try to copy file to root).
This PR wants to setup a default path for CMAKE_RUNTIME_OUTPUT_DIRECTORY when it does not exist.
* ggml-cpu: add ggml_thread_cpu_relax with Zihintpause support
Signed-off-by: Wang Yang <yangwang@iscas.ac.cn>
* cmake: enable RISC-V zihintpause extension for Spacemit builds
* readme : add ZIHINTPAUSE support for RISC-V
---------
Signed-off-by: Wang Yang <yangwang@iscas.ac.cn>
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>
* Faster tensors (llama/8)
Add fast matrix and matrix/vector multiplication.
* Use map for shader replacements instead of pair of strings
* Wasm (llama/9)
* webgpu : fix build on emscripten
* more debugging stuff
* test-backend-ops: force single thread on wasm
* fix single-thread case for init_tensor_uniform
* use jspi
* add pthread
* test: remember to set n_thread for cpu backend
* Add buffer label and enable dawn-specific toggles to turn off some checks
* Intermediate state
* Fast working f16/f32 vec4
* Working float fast mul mat
* Clean up naming of mul_mat to match logical model, start work on q mul_mat
* Setup for subgroup matrix mat mul
* Basic working subgroup matrix
* Working subgroup matrix tiling
* Handle weirder sg matrix sizes (but still % sg matrix size)
* Working start to gemv
* working f16 accumulation with shared memory staging
* Print out available subgroup matrix configurations
* Vectorize dst stores for sg matrix shader
* Gemv working scalar
* Minor set_rows optimization (llama/4)
* updated optimization, fixed errors
* non vectorized version now dispatches one thread per element
* Simplify
* Change logic for set_rows pipelines
---------
Co-authored-by: Neha Abbas <nehaabbas@macbookpro.lan>
Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local>
Co-authored-by: Reese Levine <reeselevine1@gmail.com>
* Comment on dawn toggles
* Working subgroup matrix code for (semi)generic sizes
* Remove some comments
* Cleanup code
* Update dawn version and move to portable subgroup size
* Try to fix new dawn release
* Update subgroup size comment
* Only check for subgroup matrix configs if they are supported
* Add toggles for subgroup matrix/f16 support on nvidia+vulkan
* Make row/col naming consistent
* Refactor shared memory loading
* Move sg matrix stores to correct file
* Working q4_0
* Formatting
* Work with emscripten builds
* Fix test-backend-ops emscripten for f16/quantized types
* Use emscripten memory64 to support get_memory
* Add build flags and try ci
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* Remove extra whitespace
* Move wasm single-thread logic out of test-backend-ops for cpu backend
* Disable multiple threads for emscripten single-thread builds in ggml_graph_plan
* Fix .gitignore
* Add memory64 option and remove unneeded macros for setting threads to 1
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* ggml : add GGML_SCHED_NO_REALLOC option to disable reallocations in ggml_backend_sched
Enabled in ggml-ci for testing.
* llama : update worst-case graph for unified cache
* ci : disable op offload in some tests
* fix spelling
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
This commit removes the "-dirty" suffix from the GGML version string.
The motivation for this change is to ensure that the version string
works with different ways of checking out ggml and using it in projects.
By removing the dirty flag from the version string, we avoid potential
artifacts like shared libraries getting a -dirty suffix in their names.
Instead, if the project is built from a dirty git state, the dirty flag
will be appended to the commit hash in the GGML_BUILD_COMMIT variable.
This will enable users to still identify that the build was made from
from a modified/dirty state even though the version might match a "real"
version.
For example, the commit can be produces as follows:
```c++
printf("commit: %s\n", ggml_commit());
```
Which would print the following for a dirty build:
```console
commit: 781baf2a-dirty
```
Refs: https://github.com/ggml-org/ggml/pull/1363#issuecomment-3569691546
* model: add support for extra bufs for all devices
* hexagon: add experimental ggml-hexagon backend for the Hexagon NPU
This commit introduces a new experimental backend `ggml-hexagon` with support for the Hexagon NPU.
Highlights:
- Supports Hexagon versions: v73, v75, v79, and v81
- Targets Android devices based on Snapdragon SoCs: Gen3, 8-Elite, and 8-Elite Gen5
- Supports Q4_0, Q8_0, MXFP4, and FP32 data types
- Implements core LLM ops: MUL_MAT/MUL_MAT_ID, ADD/SUB/MUL/ADD_ID, RMS_NORM, ROPE, GLU/SWIGLU, SOFTMAX
**Note:** This backend is experimental and may exhibit instability or limited performance across supported devices.
It is intended for early testing and feedback from llama.cpp/ggml developer and user community.
Co-Authored-By: Rajdeep Ganguly <rganguly@qti.qualcomm.com>
Co-Authored-By: Todor Boinovski <todorb@qti.qualcomm.com>
* hexagon: fix format checker errors
* hexagon: update readme and cmake presets
* ci: add android-ndk-build jobs that build plain ARM64 and Snapdragon versions
* hexagon: add simple graph optimizer for stacking MUL_MAT ops with the same input
* hexagon: move ADB helper scripts into scripts/snapdragon/adb
* hexagon: replace all f/printfs with GGML_LOG_...
* readme: add hexagon to the list supported backends
* hexagon: stack malmuts with quantized inputs only
* hexagon: add TODO for fixing issues in hexagon_graph_optimize
* hexagon: update to hex-sdk 6.4.0 and add scripts for running on QDC
* scripts: fix lint errors
* scripts: update qdc pytest script to make linter happy
* hexagon: add reduce sum in fp32
* hexagon: reduce number of vector stores in matmul output
* hexagon: remove the need for vdelta in reduce-multiply-x8
* hexagon: consistent use of reduce_sum_fp32 for row_sums
* hexagon: some more matmul optimizations and comments
Optimize cases where tensor dims are not multiple of 1024 (e.g in Qwen models).
We've handled those cases already but at a higher overhead.
* hexagon: update cmake presets
* hexagon: add OPMASK support for run-bench.sh wrapper
* hexagon: update to use GGML_BACKEND_API
* hexagon: remove unused logic for setting tensor flags for the views
* hexagon: add asserts to set/get_tensor to make sure we handle complete tensors
Same asserts as the CPU backend.
* hexagon: use cpy_tensor slow path for non-host buffers
* hexagon: error checks in the buffer allocator
* cmake: move include(extProj) under ggml-hexagon
* hexagon: don't forget to delete the backend on free
* hexagon: set/get_tensor size assert apply only to quantized tensors
* hexagon: reintroduce HEX_VERBOSE wrapper for GGML_LOG_DEBUG for now
GGML_LOG_DEBUG is always enabled for test-backend-ops and the output gets in the way.
Ideally we need a bit more finer log levels.
* docs: typos in hexagon developer docs (libggm-...)
* hexagon: overhaul error handling in the session/device allocation
this should handle all failure paths in the session allocation.
* hexagon: update cmake presets to enable fp16 vectors
* hexagon: remove unused time_usec function
* hexagon: don't forget to release buffer contexts
* hexagon: fixed indents in hvx-utils (missed clang-format auto-format failure)
* hexagon: remove custom can_repeat function and use ggml_can_repeat
---------
Co-authored-by: Rajdeep Ganguly <rganguly@qti.qualcomm.com>
Co-authored-by: Todor Boinovski <todorb@qti.qualcomm.com>
Georgi Gerganov
Max Krasnyansky
Neo Zhang
Katostrofik
Steve Lhomme
uvos
texasich
Johannes Gäßler
Taimur Ahmad
Zijun Yu
Adrien Gallouët
Kevin Pouget
Boian Berberov
Alfred
Congcong Cai
ixgbe
Vishal Singh
Reese Levine
xiaobing318
Diego Devesa
Daniel Bevenius
Aleksei Nikiforov