Limit the shared prefix to the first required output position of the parent and each child task.
This keeps decision output embeddings available in the final batch when questions share option tokens.
Fixes#30482
Signed-off-by: aiqubits <aiqubits@hotmail.com>
* chat : accept the think block in the qwen3-coder response format with reasoning_format none
Assisted-by: Claude Opus 5.5
* chat : define both qwen3-coder think block variants up front
This commit extracts all apple jobs except for macos-latest-arm64 into a
new daily workflow.
The motivation for this is to reducd the load on the macos runners.
The duplicated output looks up token_embd in its buffer context, so
token_embd must exist first. With the old order a load where input and
output share a buffer type (CPU-only, -ngl 0) held the table twice:
gemma-4-E2B-it Q4_K_M, no mmap, CPU model buffer 3163.73 -> 2947.73 MiB.
* memory: refuse rollbacks past the snapshots of the last ubatch
The recurrent memory and the DSV4 cache accepted any rollback up to
n_rs_seq, whether or not the last ubatch wrote that snapshot. Each now
records the deepest rollback its last ubatch guarantees and seq_rm
refuses anything past it:
- recurrent: n - 1 for a ubatch of n <= n_rs_seq + 1 tokens, 0 for a
longer one, which only guarantees the final state
- DSV4: min(n, n_rs_seq), the snapshots are rebuilt from the state
before the ubatch
The dead rollback block of the rollback test becomes the assertion it
asked for, and rollback_bound checks the exact bound of each memory
after ubatches of 1, 2, n_rs_seq + 1 and n_rs_seq + 2 tokens.
* tests: address review from ggerganov
Detect architectures with llama_arch_from_string in comp_padding and
arch_reserves_single_seq, as in rollback_bound.
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>
* ui: land at the bottom of a conversation every time it opens
Opening a long conversation often stopped short of its end:
- the press grace counted from page load (pressedAt = 0), so a chat that
landed within 400 ms of a refresh did not follow the rows realizing
after the jump
- the click that opened a chat (sidebar row, tab) started the same grace,
so navigating by click had the same problem
- any scroll event away from the bottom stopped following, including the
ones the browser fires itself while the chat settles
- the chat screen outlives its routes, so going back to the last landed
conversation through the new-chat screen skipped the landing entirely
scrollToBottom() now clears the grace, only an upward scroll away from the
bottom stops following, and every navigation lands anew. Navigations into
a chat also disable SvelteKit's scroll reset, which moved the previous
conversation to its top for a few frames.
Where overflow-anchor is supported, a 1px pin after the messages is the
only scroll anchor while following, so the browser keeps the bottom in
place itself; WebKit keeps the resize observer.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: let nested agentic content scrollers chain to the chat
Reasoning, shell output and code blocks used overscroll-behavior:
contain, so the wheel stopped dead at their edges instead of
scrolling the conversation.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: stop following the stream as soon as the reader scrolls up
A trackpad or smooth wheel scrolls in steps smaller than the at-bottom
threshold, so each step up was re-stuck and snapped back on the next
growth while a response streamed. Any upward move that is not content
shrinking now releases at once.
The attach effect also read the hook's own state, so every flip
re-attached and re-measured, which undid such a release.
Assisted-by: Claude (Opus 5.5 Medium)
ChatFormActions gated the mic and stop buttons on a `canSubmit` prop that
ChatForm never passed, so it was always false and typing text did not swap
them for the submit button. Use `canSend` instead and also hide the
skip-reasoning button while there is content to send.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: reveal the collapsed sidebar from its logo and toggle it there
The collapsed sidebar now shows only the logo; hovering it reveals the
icon strip until the pointer leaves. The expanded sidebar drops its
top-right collapse button and the logo collapses it on hover instead.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: fade the peeked sidebar in at full width and keep the toggle slide
Hovering the logo of the collapsed sidebar now fades in the fully open
sidebar, floating over the page at its full width, instead of resizing
it. Opening or closing with the toggle still resizes and slides: the
sidebar docks to strip width while it slides, then floats out of sight,
and it stays hidden under the pointer after a toggle close. The unused
icon-only strip is removed from the sidebar actions.
Assisted-by: Claude (Opus 5.5 Medium)
* ui: stack the sidebar above the chat tabs on desktop
The peeked sidebar floats over the page, but the chat tab bar shared its
z-index and came later in the page, so the tabs were drawn on top of it.
Assisted-by: Claude (Opus 5.5 Medium)
MCP server groups in the settings Tools tab now get a group checkbox like the
built-in, server and browser groups, toggling the server-scoped tools key.
Indeterminate checkboxes now use the checked state colors.
Assisted-by: Claude (Opus 5.5 Medium)
Ask the browser to keep IndexedDB and localStorage data under storage
pressure, so conversations and settings are less likely to be evicted.
Assisted-by: Claude (Opus 5.5 Medium)
* gguf-py : add Limite arch
Assisted-by: Claude Opus 5.5
Signed-off-by: Giuseppe Scrivano <giuseppe@scrivano.org>
* convert : add Limite support
Assisted-by: Claude Opus 5.5
Signed-off-by: Giuseppe Scrivano <giuseppe@scrivano.org>
* llama : add Limite support
Assisted-by: Claude Opus 5.5
Signed-off-by: Giuseppe Scrivano <giuseppe@scrivano.org>
---------
Signed-off-by: Giuseppe Scrivano <giuseppe@scrivano.org>
* 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.
Requested in the #31125 review. The cases duplicate coverage that is
already in the file and add CI time. Test-only, nothing changes in the
kernel, the docs or the CMake option.
Assisted-by: Claude
* 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.