Konrad Moren [Mon, 3 Aug 2026 12:26:09 +0000 (14:26 +0200)]
CUDA: Add backend sampler for penalties sampler (#25262)
* sampling: enhance penalty handling in common_sampler_init
- Set default value for penalty_last_n based on model context if not specified.
- Ensure penalty_last_n and n_prev are non-negative.
- Update llama_sampler_penalties structure to inherit from llama_sampler_backend and add backend input handling for penalties.
- Implement backend initialization and application logic for penalties, including frequency and presence adjustments.
* tests: add backend penalties sampling tests and utility functions
- Introduced `accept_prompt` and `unique_prompt_tokens` functions to handle prompt acceptance and token uniqueness.
- Implemented `compare_penalties_logits` to compare logits from backend and CPU samplers with penalties.
- Added `test_backend_penalties_sampling` to validate backend penalties with various configurations.
- Enhanced the test suite for better coverage of penalty handling in sampling.
* sampling: add support for top-k penalties in backend sampling
* sampling: add fix to ensure stable numerical results. Preserve masked logits as -Inf and no longer generate NaN.
* sampling: enhance penalty comparison tests with masking penalties logic
* add comments on padding
* sampling: add comments on modifications
* add the unit test to cover masked-out token as -INF
* validate repeat penalty to ensure it is finite and greater than 0; add tests for invalid values
* refactor: test functions to share logic and be less verbose
* add test to cover case where previously penalized token is not part of candidates
* remove comments
* remove redundant penalty_last_n initialization and validation in common_sampler_init
* add support for penalties in sampler chain with configurable positions
* add validation for penalty parameters and enhance tests for non-finite values
* add context parameter to common_sampler_init and set default for penalty_last_n
* add llama_n_ctx parameter to common_sampler_init for improved sampler initialization
* replace penalty_last_n x n_candidates comparison matrix with a vocabulary-sized count tensor
* add tests for backend penalties sampling without filler entries , token_count.size() == n_active == n_max == 64
* add test for backend penalties sampling after top-p with large history window
Thiago Padilha [Mon, 3 Aug 2026 04:33:37 +0000 (01:33 -0300)]
metal: implement DSv4 Lightning Indexer (#25893)
* metal: implement F16 Lightning Indexer
- Implement GGML_OP_LIGHTNING_INDEXER for 128-dimensional, 64-head inputs
with F32 queries and weights plus F16 keys and masks.
- Add tiled and tail kernels and test KV lengths around 8- and 64-element
boundaries.
Assisted-by: Codex
* metal: stage Lightning Indexer K tiles
- Stage and dequantize K in F16 threadgroup memory before simdgroup matrix loads.
- Zero-fill partial tiles and guard stores so all KV segments use the same numerical path.
- Support F32, F16, BF16, Q4_0, Q4_1, Q5_0, Q5_1, and Q8_0 K caches.
llama-bench (--mmap 1, -fa on, -p 512, -n 128; d=0/10k/20k):
- Implement GGML_OP_DSV4_HC_COMB, GGML_OP_DSV4_HC_PRE, and
GGML_OP_DSV4_HC_POST with SIMDgroup register and shuffle optimized kernels.
- Add Metal dispatch and support plumbing and test the production Sinkhorn
iteration count and embedding width.
Pascal [Sun, 2 Aug 2026 17:25:27 +0000 (19:25 +0200)]
common: support the DSpark sidecar resolution (#26458)
The dspark- files resolve like the other speculative sidecars: the
-hfd tag applies to them, a requested sidecar resolves without a full
model at the tag, and an explicit -md selection disables the discovery.
When no type is requested, dspark outranks dflash in the auto-selection
since its sidecar carries the extra Markov head.
akleine [Sun, 2 Aug 2026 13:43:00 +0000 (15:43 +0200)]
opencl: bugfix increment ref_count in ggml_backend_opencl_init() (#26162)
Incrementing `ref_count` at the beginning is important later
in the `free()` method of the `ggml_backend_opencl_context` at program end.
If we do not increment the `ref_count`, the result would be -1 here,
and consequently, the profiling data would not be flushed and written.
( #ifdef GGML_OPENCL_PROFILING )
Nico [Sat, 1 Aug 2026 16:03:32 +0000 (12:03 -0400)]
cli : persist reasoning_content in chat history (#26362)
* cli : persist reasoning_content in chat history
llama-cli collected reasoning from the stream for display but only
stored assistant content in messages, so --reasoning-preserve could
not re-inject prior thoughts on later turns.
* fattn-mkl: fix interleaved dst layout in normalize kernel
- Fix mkl_fa_normalize_head: use interleaved dst layout
((query * n_q_heads + head) * DV) matching TILE's
flash_attn_combine_results. Previously used dense head-major
layout which wrote head outputs to wrong addresses, corrupting
attention for all models except Qwen3.6-27B (where GQA=6 heads
were sparse enough to avoid visible overlap).
- Remove 7 redundant stream->wait() calls — SYCL in-order queue
already serializes pure SYCL kernel dependencies. Retain only
the 4 MKL GEMM ↔ SYCL handshake barriers (oneMKL GEMM uses its
own internal queue that does not respect SYCL in-order).
- Remove unused dst_row_stride, diagnostic clutter, and dead
K/V hex dump (fa_diag block in fattn-mkl.cpp).
- Add MKL_FA_DISABLE=1 env var for A/B testing.
- Add FA-DISP watchdog (MKL_FA_DEBUG=1) and FA-DIAG output
fingerprint (MKL_FA_DIAG=1) in fattn.cpp.
Co-Authored-By: Claude Code on DeepSeek-v4-Pro
* Thank you for the review feedback: rename env vars, use GGML_LOG_INFO, document in SYCL.md
Completed the following:
- Rename MKL_FA_DISABLE → GGML_SYCL_ENABLE_MKL_FA (inverted: 0 to disable)
- Rename MKL_FA_DEBUG → GGML_SYCL_MKL_FA_DEBUG
- Rename MKL_FA_DIAG → GGML_SYCL_MKL_FA_DIAG
- Replace fprintf(stderr, ...) / fflush(stderr) with GGML_LOG_INFO() macro
- Document all three env vars in docs/backend/SYCL.md under Runtime
- Add comment explaining MKL FA activation trigger (flash-attn + quantized
KV cache + batch-size >= 1024 + n_kv >= 1024)
Resolves review feedback from arthw.
Again, thank you!!!
Co-Authored-By: Claude Code on DeepSeek-v4-Pro
* Thank you for the review feedback round 2: use ggml_sycl_get_env, remove dup waits, gate perf macros
- Replace raw getenv() with ggml_sycl_get_env() in all 4 env-var checks
(fattn.cpp: GGML_SYCL_ENABLE_MKL_FA, GGML_SYCL_MKL_FA_DEBUG,
GGML_SYCL_MKL_FA_DIAG; fattn-mkl.cpp: GGML_SYCL_MKL_FA_DEBUG)
- Remove duplicated stream->wait() before ev.wait_and_throw() in GEMM
KQ and GEMM VKQ — ev.wait_and_throw() already waits for completion
- Gate MKL_ACCUM macro behind do_print so timing accumulators are
no-ops in normal operation
- Remove redundant MIT/Intel copyright header from fattn-mkl.cpp
- Remove unused #include <cfloat>
- Expand SYCL.md MKL FA docs with step-by-step activation trigger
and example llama-cli command
Again, thank you!!!
Co-Authored-By: Claude Code on DeepSeek-v4-Pro
* fattn-mkl: enable MKL FA for all KV cache types
Remove the quantized-only restriction on MKL activation — the MKL
kernel converts any non-F16 K/V to F16 via to_fp16_sycl before GEMM,
so F16 (default), BF16, and F32 caches all benefit from XMX hardware
acceleration. The type restriction was an unnecessary gate.
Before (F16/BF16 default cache + FA on at 32K prefill): ~356 t/s (TILE path)
After: ~670 t/s (MKL path, matching quantized-cache baseline)
Minimal change: two conditions removed, one comment updated in fattn.cpp.
No kernel or conversion code changes — the dequant pipeline already
covers all types.
* fattn-mkl: rename mkl_disable -> mkl_enable for clarity
* fattn-mkl: refine MKL FA dispatch gates
Three changes:
1. Remove quantized-only restriction - MKL FA activates for all
KV cache types (F16 default, BF16, F32, quantized). The MKL
kernel converts non-F16 K/V via to_fp16_sycl before GEMM.
2. Rename mkl_disable -> mkl_enable to match env var
(GGML_SYCL_ENABLE_MKL_FA).
3. Replace batch-size threshold with Q->ne[1] >= 32 gate.
Keeps TG (Q=1) and MTP drafts (Q=3-8) on VEC path where
fused kernel beats MKL launch overhead. Routes all
multi-token prefill through XMX-accelerated GEMM.
Production data confirms Q patterns: 1-8 TG, 32-127 cache reuse,
128+ full reprocess. At 32K F16/BF16 FA-on: 356 -> 670 t/s.
1. Always copy F16 K/V to dense row-major buffers before MKL GEMM.
Previously F16 was read in-place with raw tensor strides. During
multi-turn conversations, the accumulated KV cache had different
stride properties than a fresh prefill, producing corrupted outputs.
Now dense F16 gets a fast memcpy; interleaved (Gemma) gets a strided
copy kernel. This matches what the quantized paths already did through
to_fp16_sycl.
2. Gate MKL FA on unsupported op params (max_bias, logit_softcap, batch
dim mismatch) and pathological F16 strides (nb[1] not a multiple of
ne[0]*2). These conditions would previously crash inside the MKL
kernel. Pathological strides (test-only) and ALiBi/softcap fall
through to TILE/VEC which handle them correctly.
The stride check uses modulo rather than equality, so both dense
(nb1 == ne0*2) and interleaved (nb1 == H * ne0*2) pass — all real
models use these layouts. Only test cases with overlapping rows
(nb1=32 or nb1=75 for ne0=40) are blocked.
Thanks to hmscider for the oneDNN FA PR (#25222) which surfaced the
same insight: always normalize inputs to contiguous F16 before GEMM.
Co-Authored-By: Claude Code using DeepSeek-V4-Pro <redacted>
* fattn-mkl: fix quant+GQA KV strides, tighten MKL gate, add K>=1024 tests
Adding K>=1024 flash-attn test cases surfaced several MKL bugs:
- Quant K/V with a padded seq-view (real KV cache) used the wrong
strides in the dequant path... only the true Gemma interleave
layout should reconstruct strides. nb[2] vs ne[1]*nb[1]
- Gate was firing on shapes the kernel doesn't handle: head_dim < 64
or not a multiple of 64, MHA, attention sinks, and
bf16 decode... fell through to vec which no bf16 case.
Gate MKL to the validated envelope: gqa>=2, head_dim 64 through 512
(has to be a multiple of 64) with matching K/V head size, mask,
no sinks/alibi/softcap... everything else falls back to tile.
Covers Qwen Dense/MoE and Gemma4 Dense/MoE
Ran test-backend-ops -o FLASH_ATTN_EXT: 3641/3641 pass.
Perplexity unchanged... 6.7267 MKL vs 6.7290 stock using
Qwen 27b q5_k_xl
Co-authored-by: Neo Zhang <redacted>
* fattn-mkl: bound attention scratch so it doesn't grow with batch or context... also dropped the bf16 comment in fattn.cpp per arthw review.
* docs : center badges, remove Hot topics, extract sections, remove tools
- Use <div align="center"> for GitHub-compatible centering
- Add dev branches and compile times links
- Add lib llama API and llama-server REST API links
- Remove Hot topics section
- Remove Recent API changes section
- Extract XCFramework section into docs/xcframework.md
- Extract Completions section into docs/completions.md
- Extract Obtaining and quantizing models into docs/models.md
- Remove tools usage sections (llama-cli, llama-server, etc.)
- Move Contributing section to the end
ggml-cuda : disable MMQ on devices with less than 48 KiB shared memory (#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:
Georgi Gerganov [Wed, 29 Jul 2026 11:59:44 +0000 (14:59 +0300)]
server : add trace logging for slot similarity checking (#26271)
Adds trace logging in server-context.cpp for slot similarity checking
during prompt cache slot selection, including skip reasons and similarity
calculation details.
model: add NextN/MTP speculative decoding support for GLM_DSA (GLM-5.2) (#25980)
* 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 <redacted>
* 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 <redacted>
---------
Hongqiang Wang [Tue, 28 Jul 2026 18:04:42 +0000 (11:04 -0700)]
opencl: skip the Adreno KQ/KQV image kernels for multi-stream batches (#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).
Daniel Bevenius [Tue, 28 Jul 2026 15:20:25 +0000 (17:20 +0200)]
mtmd : add Nemotron 3 Nano Omni support (parakeet) (#22520)
* 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.
* 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 : 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 <redacted> Co-authored-by: Xuan Son Nguyen <redacted>
ggml-cuda: add chunked SSD matmul for Mamba-2 prefill acceleration (#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.
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:
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 <redacted>
* 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 <redacted>
* 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 <redacted>
* 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.
sycl: fix use-after-return of the SDPA scale in the oneDNN flash-attention path (#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 <redacted>
* 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.
common: fix explicit -md precedence over draft sidecar resolution (#26165)
* common: fix explicit -md precedence over draft sidecar resolution
Follow-up of #25955, an explicit --model-draft file given with -hfd
was silently overridden by the sidecar resolution of the draft repo,
and its path was never resolved to a local file.
An explicit draft file selection now disables the sidecar resolution,
so the manual CLI configuration wins over the automatic one.
* common: apply the -hfd tag to the sidecar resolution
The sidecar selection was anchored on the primary of the draft plan,
so a tag without a matching full model aborted the whole plan, and
the sidecar quant silently followed the default model pick.
The tag now anchors the sidecar directly: exact tag match first, then
closest quant to the tag, and a requested sidecar resolves even when
no full model matches the tag. A wired draft sidecar also counts as
an explicit draft, so the main plan no longer downloads a second one.
* common: promote speculative load logs from trace to info
Show the loaded draft model and the MTP draft context at the default
verbosity, for consistency with the mmproj and primary logs.
Co-authored-by: Georgi Gerganov <redacted>
---------
overview / pkg / ggml / sources / llama.cpp / log