* hexagon : fix DMA transfers with strides > 24 bits
On v75+ the 2D DMA descriptor stores row_size and strides in 24 bits.
dma_queue_push sent oversized strides to fallback_2d, which only splits
nrows, so strides of 16 MB or more were truncated - wrong addresses.
Fix will fall back to one transfer per row in that case.
Also added two CONT tests with source strides over 16 MB.
* hex-dma: refactor the fallback paths to remove duplication
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
gfx117X is more like RDNA4, with 128b WMMA. But they're not identical. Adjust the defines accordingly to have a new RDNA4_M and use it where it makes sense.
Assisted-by: Codex
Assisted-by: GitHub Copilot
Co-authored-by: Mario Limonciello <mario.limonciello@amd.com>
* WIP OCP FP8 E4M3 support
* Fix python type-checker
* Materialize input_scales also for FP8 when converting
* Fix conversion of https://huggingface.co/RedHatAI/Qwen3.8-27B-NVFP4
1. Per row FP8 wrongly fell to NVFP4 conversion, as NVFP4 checked on
dims alone. Also check on dtype for FP8
2. Need to reshape FP8 QKV projections scales in the same way that
weights are reshaped
* Allow remaining FP8 tensors to be dequantized during conversion
* Allow tensor-strategy for FP8-only compressed-tensor ckpts
* Fix and test gguf-py behavior for FP8_E4M3
* Disallow QKV-fusion for FP8 temporarily
* Fix conversion of fused qkvz & FP8 for Qwen3-Next
* Avoid fusing FP8 for --fuse-gate-up-exps as we may have different ws
* Emit warnings when not fusing FP8 during conversion
* More FP8 conversion fixes
* Add supported for strided copies to/from FP8 in CPU
* remove unneeded scratch buffer from FP8 cpu cpy/dup path
* add explanatory comments and TODOs
* Add dedicated encoder and decoder tests for FP8
* Whitelist ggml_op_get_rows in metal instead of blacklisting
required to fix failing CI, and white-listing should be more
future-proof than adding FP8 to the black-list
https://github.com/ggml-org/llama.cpp/actions/runs/38433916498/job/115422880633?pr=27512
* Fix more conversion bugs
* Defer weight materialization during conversion
* Blacklist fp8 in metal backend instead of white-list
* Adjust indents
* Fix lint + execute tests
* Add fast CPU kernel for F8_E4M3
* Conversion fixes
* Fix remaining gaps
* Enable tied weights with shared scales
* Make mandatory-F32 rules apply to FP8 during conversion
* Fix python type checks
https://github.com/ggml-org/llama.cpp/actions/runs/39765670052/job/119780379595?pr=27512
* Fix kaleidai unit-test failures
https://github.com/ggml-org/llama.cpp/actions/runs/39765670125/job/119780383150
* Relax numeric precision constraints for pre-check
* Leave fast-FP8 CPU on RISC-V for follow-up
We already exclude arm from the fast-path
* Allow output scale tensors to be located in any gguf shard
* Move test_conversion_fp8.py test execution to behavior-check.yaml
* Apply batched suggestions from code review
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* Move mixed compressed-tensor logic into dequant_fp8 + add unit-tests for it
* Install all requirements also on Windows for behavior-check
Fixes https://github.com/ggml-org/llama.cpp/actions/runs/40535852590/job/122300758223?pr=27512
* Unmap lazy tensors before clean-up in Windows
Should fix https://github.com/ggml-org/llama.cpp/actions/runs/40539797824/job/122314782459?pr=27512
* LF line endings
https://github.com/ggml-org/llama.cpp/actions/runs/40539798168/job/122314791118?pr=27512
* Extract FP8 tests from llama-arch to its own test-fixture
See https://github.com/ggml-org/llama.cpp/pull/27512#discussion_r4400628221
* Update gguf-py/gguf/scripts/gguf_convert_endian.py
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
---------
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* metal : faster MXFP4 mat-vec kernel
Decode the fp4 values with the fp16 bit trick instead of a threadgroup lookup table, and load the packed nibbles as aligned 32-bit words. The kernel no longer uses threadgroup memory.
* metal : launch contiguous unary and binary ops with full threadgroups
The contiguous unary path and the row-broadcast binary path used one thread per threadgroup. Use up to 256 threads per threadgroup and bound-check the tail.
* metal : split GLU rows across threadgroups
A GLU row ran on one threadgroup, so single-row decode used a single threadgroup for the whole op. Split each row over threadgroups of up to 256 threads.
* hex-upscale: basic upscale ops (not fully vectorized)
* hex-upscale: optimize nearest 2x, 4x and bilinear 4x
* hex-upscale: add support for non-contig rows and dims collapsing
* hex-upscale: use vgather instead of scalar staging
* hex-upscale: move all kernel select logic to the host, review comments and cleanups
* tests: add permuted UPSCALE test to cover qwen3-vl image->text scenarios
* sycl : add DPAS flash attention kernel for quantized KV cache
Decode and small batches (1 to 31 query tokens) now run flash attention
straight on q4_0 and q8_0 KV blocks, instead of scalar dequant or
converting the whole cache to f16 per call. Q.K uses s8 x s4 and s8 x s8
DPAS on XMX, V.P stays f16. dpas.s4.s8 is inline vISA since joint_matrix
has no s4.
Covers head sizes 64 to 512, GQA up to 32 and attention sinks. Prompts
stay on the existing kernels. GGML_SYCL_FA_DPAS=0 turns it off.
B580, Llama 3.1 8B Q4_K_M, q4_0 KV, tg128: 31.3 to 54.2 t/s at 16K
context. FLASH_ATTN_EXT 4582/4584, same 2 failures as master.
Assisted-by: Claude
* sycl : 192 head dropped, Wildcat Lake added to docs
Assisted-by: Claude
* sycl : stricter checks from review
Assisted-by: Claude
* [SYCL] : rename XMX macro, add GGML_SYCL_SUPPORT_XMX option, document DPAS
Assisted-by: Claude
* sycl : widen reordered q4_0 dequantization stores
The reordered-layout Q4_0 conversion kernel processed one qs byte per
work-item with two scalar stores, strided 16 elements apart. Give it
the same treatment the Q8_0 kernel got in #30775: a flat
sycl::range<1> launch with four qs bytes (8 outputs) per work-item
inside one 32-element block and two vec<dst_t,4> stores.
Same float math and rounding as the scalar kernel; output is
bit-identical. The existing k % QK4_0 == 0 assert guarantees whole
32-element blocks, so the k/8 launch covers every row exactly. The
now-unreferenced scalar reorder kernel is removed.
Arc Pro B70, primed prefill (prefill after decode on a
reorder-installed tensor): +6.4% to +7.3% pp512/pp2048 on Q4_0
models (tested up to 27B dense); fresh-layout prefill and decode
unchanged; PPL within run noise.
* sycl : vectorize reordered q4_K dequantization stores
The reordered-layout Q4_K conversion kernel already loads its
quantized inputs through sycl::vec but writes the dequantized
results with scalar stores. Build the results in a vec<dst_t,n>
register and emit two vector stores instead. Same per-element
math; output is bit-identical.
Arc Pro B70, primed prefill (prefill after decode on a
reorder-installed tensor): +14.7% to +15.3% pp512/pp2048 on
Q4_K-dominant models (tested up to 27B dense); fresh-layout
prefill, decode and PPL unchanged.
* sycl : remap reordered q5_K dequantization to vec4 loads/stores
The reordered-layout Q5_K conversion kernel gave each thread only
four outputs fed by scalar ql/qh byte loads plus a local-memory
scale broadcast, so the stores were never its bottleneck: widening
the stores alone measured within noise on mixed-quant models and
+0.5% on a q5_K-dominant model. Re-decompose to 32 threads per
block, each consuming four consecutive ql and qh bytes and emitting
two aligned vec<dst_t,4> stores; a 4-element run shares one scale
pair, so the per-element math and the output are identical.
Arc Pro B70, primed prefill on a q5_K-dominant 27B model: +3.6%
pp512/pp2048 over the scalar kernel; fresh-layout prefill, decode
and PPL unchanged. The GET_ROWS reorder path moves to the same
32-thread decomposition.
* sycl : remap reordered q6_K dequantization to vec4 stores
The reordered-layout Q6_K conversion kernel wrote four scalars per
thread, strided 32 elements apart, each pair of neighbors sharing
neither instruction stream nor store transaction. Re-decompose the
same 64-thread work-group so each thread emits one aligned
vec<dst_t,4>: a 4-element run shares one 16-element scale group,
one ql nibble source and one qh bit pair, so the per-element math
and the output are identical to the scalar kernel.
Arc Pro B70, primed prefill: +28.7% pp512/pp2048 on a Q6_K 27B
model; fresh-layout prefill, decode and PPL unchanged.
* sycl : vectorize reordered q3_K dequantization stores
Each thread of the reordered-layout Q3_K conversion kernel already
computes four consecutive outputs; build them in a vec<dst_t,4>
register and emit one vector store instead of four scalars. Same
per-element math; output is bit-identical.
Arc Pro B70, primed prefill: +16.8% pp512/pp2048 on a Q3_K-dominant
27B model; fresh-layout prefill, decode and PPL unchanged.
* sycl : vectorize reordered iq3_xxs/iq3_s dequantization stores
Both reordered-layout IQ3 conversion kernels write eight contiguous
outputs per thread as scalar stores (four from each grid lookup).
Build two vec<dst_t,4> registers and emit two vector stores. Same
per-element math; output is bit-identical.
Arc Pro B70, primed prefill: +20.3% to +21.1% pp512/pp2048 on
IQ3_XXS- and IQ3_S-dominant 27B models; fresh-layout prefill,
decode and PPL unchanged.
* sycl : remap reordered q2_K dequantization to vec4 stores
Re-decompose the reordered-layout Q2_K conversion kernel the same
way as Q6_K: one aligned vec<dst_t,4> per thread instead of four
scalars strided 32 apart. A 4-element run shares one scale byte and
one 2-bit lane, so the per-element math and the output are
identical to the scalar kernel.
Arc Pro B70, primed prefill: +13.1% pp512/pp2048 on a Q2_K-dominant
27B model; fresh-layout prefill, decode and PPL unchanged.
* sycl : widen reordered mxfp4 dequantization stores
Give the reordered-layout MXFP4 conversion kernel the same flat
sycl::range<1> launch as Q4_0: four qs bytes (8 outputs) per
work-item inside one 32-element block and two vec<dst_t,4> stores.
Same LUT lookups and per-block scale; output is bit-identical. The
k % QK_MXFP4 == 0 assert guarantees whole blocks, so the k/8 launch
covers every row exactly. The old scalar kernel is removed.
The reordered layout for MXFP4 is installed on the MoE expert path
(reorder_qw_mxfp4_moe); dense MXFP4 tensors keep the normal layout.
Arc Pro B70, MXFP4 MoE 20B model: +57% pp2048 / +61% pp512 prefill
(primed and fresh alike), decode unchanged, PPL within the model's
run-to-run noise.
SCALE now accepts a src0 view with arbitrary row strides. translate_scale
used the raw input, and on the dynamic path a VIEW translates to its whole
source tensor, so a strided view input scaled the entire source. The model
output then had the source's size, and copying it back overran the result
buffer, which crashes test-backend-ops on the SCALE view cases.
Slice the input with process_view_input_new, as the other translators do.
Reject an in-place SCALE on a non-contiguous view: its result is copied
back densely and would land in the wrong elements.
* hexagon: IM2COL over multiple threads with output-col blocks
Assisted-by: Claude Opus 5.5
* hex-im2col: fix padding handling
* hex-im2col: loop unroll common cases
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* opencl: run the tiled q4_K lm_head through a tiled GEMM at batch > 1
* opencl: split tiled dp4a q4_k gemm into its own file and test for dp4a
Assisted-by: Codex
---------
Co-authored-by: Li He <lih@qti.qualcomm.com>
* hexagon: add conv2d op
* hex-conv2d: add double-buffered staging
* hex-conv2d: full 64-wide activation block transpose straight from staged buffer
* hex-conv2d: implicit gemm
* hex-conv2d: remove scalar path and kernel
* hex-conv2d: unify conv2d kernel params for all paths
* hex-conv2d: scale HMX result before f16 storage and then restore during unpacking to prevent overflow
* hex-conv2d: fix f16 overflow in HMX paths
* hex-conv2d: refactor to rebase to new op fusion machinery
* hex-conv2d: restructure DMA pipelines and trace scopes
* hex-conv2d: removed the HVX_WIDTH and HVX_TILED paths
* hex-conv2d: remove in-loop vector spills
* hex-conv2d: refactor kparams, vtcm-layout and dma-pipeline to match the other ops
* hex-conv2d: add basic support for CONV_2D_DW
Co-Authored-By: Pranav Dhinakar <pdhinaka@qti.qualcomm.com>
* hex-conv_2d_dw: use DMA in all paths
* hex-conv_2d_dw: support for chunking larger inputs
* hex-conv_2d_dw: remove scalar/hvx access for w tensor (replaced with sync dma copy)
* hex-conv-2d: fix warnings
* hex-conv2d-dw: vectorize all paths
* hex-conv-2d-dw: tree-based reductions and unrolling to optimize
* hex-conv-2d-dw: uninline compute loops from dma loops
* hex-conv-2d-dw: move all kernel selection into the host solver and simplify/specialize kernel dispatch
* hex-softmax: update obsolete gates to enable SOFT_MAX for mmproj
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
Co-authored-by: Pranav Dhinakar <pdhinaka@qti.qualcomm.com>
* mxfp4 sve 256 impl completed
* working with pretty good perf improvement
* review comments addressed
* review comments for SVE 256 check
* trailing white spaces removed
---------
Co-authored-by: Anubhav Srivastava <anubhav.srivastava@fujitsu.com>
* fix: clamp the A prefetch to end_k in the Vulkan coopmat1 int8 matmul, like the B prefetch
* fix: zero the A scale for k steps past end_k in the Vulkan coopmat1 int8 matmul instead of relying on the clamped block
* vulkan: allow dim 3 broadcast in MUL_MAT
supports_op rejected every MUL_MAT whose src0 and src1 differ in ne3,
although mul_mm, mul_mmq and mul_mat_vec already broadcast src0 over
dim 3 through broadcast3. Such a matmul then ran on the CPU with a
copy of the weight, and the split layout changed with the number of
streams in the ubatch.
Keep the strict check for MUL_MAT_ID only, and route the broadcast
cases away from mul_mat_vec_nc, the one path that requires equal ne3.
* tests : harden MUL_MAT with permuted dim 3 broadcast cases
Harden the MUL_MAT coverage with a weight broadcast over dim 3 against
a permuted per-stream activation, the shape of an MLA output
projection with a non unified KV cache, for a single token, several
heads and several tokens per stream.
* metal : add GGML_OP_TOP_K_MASK
Port the fused top-k mask from the CUDA backend: a grid-stride kernel fills the output with -inf, then one
threadgroup per row copies the mask values of the kept cells over it. Function constants drop the score,
blocks and extra paths that a given graph does not use.
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
* metal : clean up top-k mask kernel formatting
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
* ggml : relax scale to non-contiguous views
SCALE requires src0 to be 1D-contiguous. Relax this to dim-0
contiguity with arbitrary row strides (ggml_is_contiguous_rows),
same as the unary ops in the Metal backend.
- ggml: relax the assert in ggml_scale_impl
- cpu: address src0 rows through their byte strides (f32, bf16)
- cuda: stride-aware kernel, fastdiv row decomposition
- metal: support bf16 scale (guarded unary instantiations)
- sycl: keep the gate at contiguous, the kernel is flat f32
Add view variants to test_scale. Verified on CPU, Metal and CUDA.
Assisted-by: dsh:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (RTX 5090 32GB, DDR5 128GB)
* cuda : skip the softcap fusion on non-contiguous views
The fused softcap kernel reads src and writes dst as dense arrays,
so a view reaching SCALE -> TANH -> SCALE now that SCALE accepts
non-contiguous rows would be read wrongly. Keep such graphs on the
strided kernels until the fusion becomes stride aware.
---------
Co-authored-by: Pascal <admin@serveurperso.com>
* hexagon: add GGML_HEXAGON_VTCM_MB and GGML_HEXAGON_NTHREADS caps
Change-Id: Idffe1613c8db6f96aa842d43d1b2efaed45462a8
* hexagon: rename VTCM cap to GGML_HEXAGON_VTCM_SIZE, drop NTHREADS cap
Change-Id: If4c5b9ce9fd6b725c3561421c2295d40599d012d
* hexagon: make GGML_HEXAGON_VTCM_SIZE an override, accept K/M/G suffixes
Change-Id: Ib8eea84c22f24e6d986b37f215e504b046478c2c
Every dst row comes from src0 or src1, or for dim 0 from the row of
src0 followed by the row of src1, so a single kernel covers every
type: a dense source row is one memcpy, quantized types included,
and a strided source is copied element by element with the element
size known at compile time.
Rows are split across threads instead of i2 alone, which is often 1,
and are copied in tiles of 32: for a transposed source the elements
are walked in the outer loop, so the rows of a tile read neighboring
addresses instead of jumping a full row stride per element.
Add transposed b cases to test_concat, the layout of the delta-net
conv input, and a perf case for it.
* sycl: add set_tensor_2d for faster tensor-split loading
Pack host rows into chunks of up to 64 MiB when destination rows are
contiguous, avoiding a separate upload call for each row. Keep the
existing per-row copy path for non-contiguous destinations.
On two Arc Pro B70s with gpt-oss-120b UD-Q4_K_XL, total llama-perplexity
runtime drops from 263-268 s to 114-115 s. A full-load readback check
verified 26,784,256 rows with no mismatches.
* sycl: batch 2D tensor uploads through shared staging slots
Implement set_tensor_2d to speed up tensor-split model loading by
packing rows into the existing pinned staging slots and uploading
each batch with a single device copy.
Share staging initialization with set_tensor and wait only when
reusing a slot, allowing CPU packing and device copies to overlap.
Fall back to per-row uploads for noncontiguous device rows,
oversized rows, or unavailable staging.
---------
Co-authored-by: Jon Staples <250234834+Captain-Tripps@users.noreply.github.com>
xrip
Michael Taylor
shaofeiqiandLi He
Mark TodorovichandMax Krasnyansky
Jesse Symons
Georgi Gerganov
Slobodan JosicandMario Limonciello
Oliver SimonsandSigbjørn Skjæret
uvos
Aman Gupta
Trivikram ReddyandMax Krasnyansky
Anant Shrivastava
Ravi Panchumarthy
Marco ColomboandMax Krasnyansky
Nikhil Sharma
Bartowski
anubhavfujitsu
abhijain1204fujitsuandAnubhav Srivastava
Frosty40
anujj
Yonghye Kwon
Kris
Katara
Pascal
Ehsan Bateni
virajwad
Titaniumtown
Captain-TrippsandJon Staples
KyleHagy