* improve fa of quantized kv cache
* Fix some bugs and some comments.
* fix v type check and some comments
* Fix build error caused by rebasing
* editorconfig checking pass
* add bool cwhn = true to conv_2d test cases
* add layout check at graph building time
* extend layout checks for conv2d.cu kernel
* in CPU back-end kernel needs to be stored contiguously to prevent test failures with cwhn=1
* trim white space
* do op support check in vulkan backend
* fix CI failure and vulkan run-time assert failure by introducing new graph build-time check in ggml_backend_vk_device_supports_op
* add additional check in support_op function for Vulkan to fix run-time assert failure
* 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>
* fix: single-flight conversations store init
* refactor: remove unused legacy-migration util
* fix: make createSystemMessage transactional
* fix: delete message branches cascading on edit/regenerate
* fix: stop stamping lastModified on conversation metadata updates
* fix: count cascaded forks in bulk delete toast, bulkify deleteAll
* refactor: drop redundant conversation list respreads
* refactor: create conversation in a single write
* fix: use table constant in toggleConversationPin
* fix: keep the system message placeholder out of the edit form
* fix: keep focus in the system message editor after opening it
* fix: focus the main chat form after submitting a system message
* fix: update timestamp of the correct conversation on stream completion
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>
Adds trace logging in server-context.cpp for slot similarity checking
during prompt cache slot selection, including skip reasons and similarity
calculation details.
Assisted-by: llama.cpp:Qwen3.6-27B
* model: add NextN/MTP speculative decoding support for GLM_DSA (GLM-5.2)
Adds GLM-5.2 NextN/MTP as a --spec-type draft-mtp target: nextn tensor
loading via the qwen35moe/step35-style presence probe, a graph_mtp
builder (enorm/hnorm/eh_proj + dense MLA + sigmoid-gated MoE with
shared expert + shared head with fallbacks, _s scale tensors passed
for NVFP4), t_h_nextn extraction in the trunk graph, and MTP-context
KV setup: the draft head runs dense MLA, so the MTP context uses a
plain attention KV cache holding only the nextn layer(s) (same
pattern as the hybrid Qwen3.5 MTP context) while the main context
keeps the DSA cache, now filtered to trunk layers only.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* convert : support --mtp/--no-mtp export for GlmMoeDsaForCausalLM (GLM-5.2)
Opt GLM-5.2 into the supports_mtp_export contract (post-#25641 shape,
mirroring HYV3Model/Step35Model): --no-mtp drops the appended NextN
block (blk.78) and its nextn_predict_layers KV; --mtp keeps only the
NextN block plus shared embeddings/norm/lm_head. Default (bundled)
output is unchanged.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
---------
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* 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
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)
* mtmd : add Nemotron 3 Nano Omni support (parakeet)
This commit adds support for the subsampling and encoder part of
Nemotron Nemo 3 omni model.
The Parakeet subsampling/encoder were taken from parakeet.cpp which
is currently a pull request against whisper.cpp. I've tried to copy the
code a close as possible to hopefully enable easy patching between the
these two project later.
Refs: https://github.com/ggml-org/whisper.cpp/pull/3735
* mtmd : generate rel pos tensor in graph instead of in conversion [no ci]
This commit removes the generation of the relative positional tensor in
the model conversion script and instead computes it in the encoder
graph. This is only done for the window of positions required for the
current audio sample.
* mtmd : add clip_get_model to clip API [no ci]
This commit adds a function to get access to the clip_model. It also
removes the two functions clip_get_mel_filter_tensor, and
clip_get_window_tensor(const struct clip_ctx * ctx) which can now use
clip_get_model to access the model tensors that it needs.
* mtmd : read mel_filters and window into hparams
* mtmd : use set_input_f32 lambda [no ci]
* mtmd : add better asserts for mel_filters and hann window [no ci]
* mtmd : add missing size_t cast
* mtmd : change type of pad to size_t
* mtmd : zero initialize samples_padded
* mtmd : remove unsued ctx member from parakeet preprocessor
* mtmd : make log_mel_spectrogram_parakeet_worker_thread private static
* mtmd : sync/update parakeeet impl with latest whisper.cpp
This commit updates the parakeet code in mtmd to reflect the latest
updates to parakeet.cpp in whisper.cpp.
A follow up commit will address the currently hardcoded dw_pad and see
if we can add n_conv_kernel as a model metadata field.
* mtmd : add audio_conv_kernel_size to model conversion
This commit updates the model conversion to read the conv_kernel_size
field from the sound_config section of the models config.json file.
It then uses this field instead of the hardcoded values in parakeet.cpp.
* mtmd : cleanup [no ci]
* conversion : call super().filter_tensors [no ci]
* do not discard result of super filter_tensors
* mtmd : use build_mm instead of ggml_mul_mat
* mtmd : use build_ffn
* mtmd : move and reuse get_vector lambda
* mtmd : use build_inp_raw for parakeet
* mtmd : throw exception in get_scalar instead of assert
* mtmd : fix std::min call
* mtmt : use .c_str in throw clause in get_vector
* mtmd : check for F32 type and non-empty tensor in get_vector
The get_vector lambda is used by get_scalar but also standalone to read
in the mel_filters and the window data. Therefor we are not checking
for 1D tensors but allowing multiple dimensions. We do have a check in
get_scalar to verify the size of the vector.
* mtmd : replace hardcoded 1101 for n_tokens_real
* mtmd : assert subsampling_factor is 8
This commit adds an assert of the parakeet subsampling factor to check
that it is 8.
The motivation for this is that this model currently has three
convolutions with a stride of 2. If the underlying model updates the
subsampling factor these convolution operations will need to be updated
and this will produce and error if this occurs.
* mtmd : remove unused ggml_tensors attn_pos_w and mm_norm_w
* mtmd : remove single thread path
This commit removes the single thread path which was a left over from
the original parakeet.cpp where n_threads is configurable.
* fix some security issues
---------
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* 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-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
* spec: add DSpark speculative decoding
DSpark (DeepSpec, 2026) on top of the merged DFlash drafter. It reuses the
DFlash encoder/decoder graph, target feature extraction and KV-cache injection,
and the verify/accept path unchanged; the draft model is a new "dspark" arch
adding a low-rank Markov head (markov_w1/w2) and an optional (unused here)
confidence head. No new public APIs.
The proposal is the only change: the block is anchor-first (position 0 already
predicts the first draft) and the decoder graph applies a semi-autoregressive,
previous-token conditioned logit bias in-graph, chained per block position:
logits'(i) = logits(i) + markov_w2 . markov_w1[prev(i)]
prev(0) = the block's anchor token, prev(i>0) = argmax(logits'(i-1))
vectorized across all blocks in the batch; the anchors are fed through a
dedicated graph input (token 0 of every block). Greedy stays lossless
(verify unchanged, same as DFlash).
- new arch "dspark" (llama_model_dspark : llama_model_dflash, reuses the graph,
loads the markov/confidence tensors; shares the target's embed/lm_head).
- Qwen3DSparkModel converter.
- new spec type "draft-dspark" (common_speculative_impl_draft_dspark :
common_speculative_impl_draft_dflash, overrides draft() only: submits whole
anchor-first blocks and greedily reads back the biased logits).
* spec: read draft block size in the dflash impl
* docs: add DSpark section to speculative.md
* spec: keep dspark block size read in the dspark impl
* dspark : add TODOs for incomplete parts
- confidence head is loaded but not used yet
- confidence-scheduled prefix pruning is not implemented
- the in-graph Markov chain is greedy-only
- only Qwen3 backbones are supported for now (also noted in docs)
* spec: fold DSpark into the DFlash arch
Address review: drop LLM_ARCH_DSPARK and the dspark.block_size /
markov_rank GGUF keys. A DSpark draft now converts to a DFlash GGUF;
the Markov head tensors are detected by presence (like eagle3 d2t),
block_size is read from the existing dflash.block_size key, and the
block anchors are taken as a strided view of the decoder's token
input instead of a separate graph input.
* spec: add confidence-based draft pruning for DSpark
The DSpark confidence head predicts per-position acceptance of the
drafted block. --spec-draft-conf-min truncates the block at the first
position below the threshold (default 0 = disabled).
* fold the dspark impl into dflash, selected by spec type
* address review comments
* dspark: clean up and improve naming
* update readme
* remove trailing whitespace
* dflash: draft full n_max blocks, defer dp.n_max to the central truncation
The DSpark markov head views the draft batch as a uniform [n_seqs x block]
grid, but the per-seq dp.n_max clamp could produce blocks of different
sizes, silently corrupting the strided views and the resulting logits.
Drop the clamp and always draft the full n_max block for every sequence:
dp.n_max is already enforced by the central truncation in
common_speculative_draft(), the same way eagle3 handles it.
Co-authored-by: Zaire404 <3147879462@qq.com>
* dflash: assert the markov head block-uniformity invariant, require the conf head
With the draft batch always submitting equal-size n_max blocks, a
non-divisible token count can only mean the batch was split across
ubatches or a caller broke the layout - fail loudly instead of silently
dropping the markov bias. The block_drafts > block_size early return
stays: worst-case graph reserve passes legitimately build with
n_seq_tokens > block_size.
Also make conf_proj required when the markov head is present: the
confidence head is part of the DSpark checkpoint format, and a missing
head would otherwise leave --spec-draft-conf-min silently reading stale
embeddings instead of confidences.
Co-authored-by: Zaire404 <3147879462@qq.com>
* dspark: fold conf_min into p_min
p_min and conf_min express the same thing - the minimum predicted
survival probability for a drafted position - differing only in how the
estimate is obtained: token probability for regular drafters, the
trained confidence head for DSpark. The DSpark readback never used
p_min, so reuse it for the confidence threshold and drop the separate
--spec-draft-conf-min flag. Both defaulted to 0 (disabled), so behavior
is unchanged.
Co-authored-by: Zaire404 <3147879462@qq.com>
* dflash: note the confidence broadcast workaround
Requested in review: the ggml_repeat only adapts the [1, n_tok]
confidences to the n_embd-wide embd_nextn transport so that
llama_get_embeddings_nextn can be reused - not a placeholder.
Co-authored-by: Zaire404 <3147879462@qq.com>
* cont : clarify
[no ci]
---------
Co-authored-by: Ruixiang Wang <wangruixiang07@outlook.com>
Co-authored-by: Zaire404 <3147879462@qq.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* 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>
