* 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
* remove as unused
* add is_disabled method, tensor logits reshape, add rest review suggestions
* clarify comment
- 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.
Assisted-by: Codex
Co-authored-by: Thiago Padilha <thiago@padilha.cc>
* vulkan : add pool1d push constants and pipeline field
Declared data structures needed for POOL1D OP, which are the vk_op_pool1d_push_constants struct and pipeline_pool1d_f32 field.
* vulkan : add pool1d compute shader
Added pool1d.comp for Vulkan backend mirroring the existing pool2d shader.
* vulkan : add full GGML_OP_POOL_1D support
Added pipeline creation and op dispatch for 1D pooling in the Vulkan backend.
* vulkan : fix pool1d shader logic
Registered pool1d_f32 in vulkan-shaders-gen.cpp and fixed tensor dimension indices and avg pool scale.
* vulkan : fix pool1d end boundary crash and expand test coverage
Fixed an issue where the shader crashed when the end boundary was negative when k0 < p0. Also, added more test cases related to this fix.
* SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt processing
* 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.
Tested: Gemma-4-26B, Gemma-4-31B, Qwen3.6-27B, Qwen3.6-35B-A3B
Perf (B70/Battlemage, 32K, q8_0 KV):
Gemma-4-26B: 1473 t/s MKL vs 746 TILE (1.97x)
Qwen3.6-27B: 609 t/s MKL vs 330 TILE (1.85x)
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.
* ggml-sycl: fix F16 cache + MKL FA multi-turn corruption; add gate guards
Two changes:
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 <noreply@anthropic.com>
* 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
* Update ggml/src/ggml-sycl/fattn.cpp
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* Update ggml/src/ggml-sycl/fattn.cpp
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* Update ggml/src/ggml-sycl/fattn.cpp
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* 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.
* Update ggml/src/ggml-sycl/fattn-mkl.cpp
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* Update ggml/src/ggml-sycl/fattn-mkl.cpp
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* apply arthw suggestions: enum for dequant modes, macro for wg_size, env-var one-liners
---------
Co-authored-by: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com>
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* 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>
* 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
* 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
* 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>
* add common/subproc.h|cpp
* add compile flag LLAMA_SUBPROCESS
* disabled by default on android and ios
* test-jinja: use common subproc
* mtmd: disable video if subproc is not set
* disable subproc on wasm
* make is_created atomic
* migrate server-mcp
* Add preliminary MiniMax-M3 support
Text-only port that re-uses existing components: MiniMax-M2 style GQA with
per-head QK-norm and partial rotary, DeepSeek-V3 style leading-dense and
routed/shared experts, and swigluoai activation. Sparse attention is not
yet supported (dense fallback); vision tower and MTP heads are dropped.
* MiniMax-M3 vision tower (mmproj + clip graph)
* Delete m3_vision_ref.py
* Update clip.cpp
* MSA
* Update constants.py
* Update minimax.py
* Cache creation. Working withotu flash attention
* Added flash attention for sparse layers
* Decomposed slow cpu OP into GPU + CPU ops. Massive speedup over long ctx
* Rewrote indexer op to be cuda native. Modified flash attention to match per group block picking
* Implement sparse attention calc out of stock ops.
* Fix a cache allocation and cont issue
* Fixed -fa auto crash, flagged debug spots
* Delete vocab.json
* Delete model.safetensors.index.json
* Delete generation_config.json
* Delete Minimax directory
* Handled multi stream case to fall back on Dense Attention
* Development scaffolding cleanup. No functional change to the decode or
4-way paths. Full debug harness remains at <8136a9c68ed7a5eb009aa67bba3fda8062f4648f> for reproducing the
selection-parity validation.
* Remove redundant comment from minimax-m3.cpp
* Changed 3 Gelu Ops for vision into Gelu_erf ops
* Assert that n_kv is multiple of 128
* Rename MSA index tensors to indexer convention
Note: All GGUFs generated before this change will need to be regenerated.
* Fix incorrect Assert
* Review driven changes (#3)
* Remove comment from conversion minimax.py
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
* Remove whitespaces from constants.py
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
* Tighten comment in minimax.py
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
* inherit MiniMax-M3 from MiniMax-M2
* drop dead text_config fallbacks
* Add indexer writer methods
* Reuse LLM_FFN_SWIGLU_OAI_MOE
* Remove duplicate indexer setters, add only block_size/local_blocks, follow value naming convention
* Fix conversion error /gguf_writer.py
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
* Update gguf-py/gguf/gguf_writer.py
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
* Update gguf-py/gguf/tensor_mapping.py
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
* Update conversion/minimax.py
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
* Update conversion/minimax.py
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
* Remove whitespace in src/llama-kv-cache.cpp
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
* Remove Whitespace in Update src/llama-model.h
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
* Remove whitespace in src/llama-hparams.h
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
* remove multimodal code upon maintainer request. Will be made as a separate PR
* Whitespace clean in tensor_mapping.py
* Log cache size on launch, block ctx shift, support prompt caching
Log indexer cache size on launch
Disallow ctx shift
Support prompt caching
* Update minimax-m3.cpp
* Optimize implementation, add multi stream support.
Fully rewrote minimax-m3.cpp for speed and buffer size gains:
Unified the 4-way + decode, 1 FA call per layer instead of 4, with the groups mapped onto ne[3]
Custom CPU op now emits block-level mask, expanded on GPU, which causes CPU to GPU transfer to shrinks at prefill
Decode: ~25 nodes/layer vs ~50, no per-group concats/conts
Unified selection semantics, so both regimes rank bs + local bias (position-anchored local force), which means prefill/decode can no longer disagree on selection
can_reuse on the MSA bias input. Graph reuse at decode restored (was rebuilding the full graph every token)
In-place mask adds, shrinking compute buffer ~6.8 to ~4.2 GiB at ub2048/62k
Multi-stream: MSA now runs with -np N when kv_unified=false. Decode stays batched across streams (still 1 FA call), prefill loops per stream. dense fallback only for --kv-unified + multi-seq
Measured effect on expert offload bound setup: decode 6.2(4WAY)–7.15(MSA_decode) -> 7.7~7.8 t/s, flat from 5k to 60k+. prefill around 10% faster. buffer about 20% smaller, multi-user support.
* set default cache type to F32
* Fix potential DSA double indexer cache allocation bug, only allocate in-cache k_idx for archs that opt in
* remove F16 downcasts in MSA attention, force F32 indexer score accum
* Add Minimax eos to llama vocab
* Guard edge case where idx cache can become stale after a tail trim
* Update llama-kv-cache.h
* Update llama-kv-cache.cpp
* Update llama-kv-cache.cpp
* Update llama-kv-cache.h
* Update llama-kv-cache.cpp
* Review driven changes
* style fix
* indexer hparams are required
* fix tests
* fix lint
---------
Co-authored-by: Daniel Han <danielhanchen@gmail.com>
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* common : extract trie/ac to a separate file
* common : support multiple token sequences in the reasoning budget sampler
* common/trie : return matched word index
* common/trie : rename "word" to "pattern"
* common/reasoning-budget : expose matched end sequence
* common/sampling : replay end sequence when reasoning budget is done
* cont : update to use multiple end sequences
* cont : clean up
* chat: fix DS4 template to explicitly follow reference behavior
* Support DeepSeekv4 flag (`drop_reasoning`).
* fix: hook DS3.2 parser for DS4 as well
* fix: add tool result reordering
* fix: post-merge
* 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
* 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
* 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>
* model: add Hy3 (hy_v3) architecture support
Adds Tencent Hunyuan 3 (HF architecture HYV3ForCausalLM, GGUF arch
hy_v3): a MoE decoder stack with per-head Q/K RMSNorm, a sigmoid
router with expert selection bias, an always-active ungated shared
expert, and leading dense block(s) (first_k_dense_replace).
The base implementation is ported from charlie12345's fork
(https://github.com/charlie12345/ROCmFPX, src/models/hyv3.cpp),
adapted to current mainline APIs (hparams.n_layer(), build_qkv,
build_moe_ffn with fused gate_up + scale tensors, output_s).
Note: blk.N.exp_probs_b is stored without a .bias suffix for
compatibility with existing hy_v3 GGUFs produced by that fork.
Co-Authored-By: charlie12345 <charlie12345@users.noreply.github.com>
Co-authored-by: Piotr Wilkin <ilintar@gmail.com>
Assisted-by: Claude Fable 5
* tests: Harmonize the use of private ggml includes
* tests: In test-backend-ops, use quoted includes
As with all other tests. This is to ensure that the build uses shipped
headers over possibly system-installed ones.
* chat : fix reasoning leak with force-opened bare <think> templates
The reasoning start tag inferred from prior turns can carry trailing
whitespace (e.g. <think>\n) while a force-open template prefills a bare
<think>. Trim the tag used for the prefix split so the bare prefill is
matched instead of being swallowed into content.
* chat : fix Nemotron Nano v2 regression
---------
Co-authored-by: Alde Rojas <hello@alde.dev>
* server: honour per-request reasoning_budget_tokens in chat completions
The reasoning-budget block in oaicompat_chat_params_parse read only the
server-level default (opt.reasoning_budget, typically -1) and the
Anthropic-style alias thinking_budget_tokens, but never the canonical
reasoning_budget_tokens field from the request body. Because the key
was then written into llama_params before the generic body-copy loop
ran, the copy loop found the key already present and silently skipped
the caller-supplied value. Any per-request override (e.g. 0 to
suppress thinking entirely) was therefore discarded.
Fix: read reasoning_budget_tokens from the request body first, so the
value that reaches the sampling layer is the one the caller intended.
Add a unit test in test-chat.cpp that exercises this path via
oaicompat_chat_params_parse with a Qwen3 template (which the autoparser
detects as a thinking-capable model) and asserts the returned
llama_params carries reasoning_budget_tokens == 0.
* server: honour per-request reasoning_budget_message in chat completions
The reasoning-budget block in oaicompat_chat_params_parse wrote
reasoning_budget_message into llama_params straight from the server-level
default (opt.reasoning_budget_message) and never read the canonical
reasoning_budget_message field from the request body. Because the key
was written before the generic body-copy loop ran, that loop found the
key already present and silently skipped the caller-supplied value. Any
per-request override of the message injected before the end tag when the
budget is exhausted was therefore discarded, even though server-task.cpp
already reads reasoning_budget_message from that data.
This mirrors the reasoning_budget_tokens bug fixed in the previous commit.
Fix: read reasoning_budget_message from the request body first, falling
back to the server default, so the value that reaches the sampling layer
is the one the caller intended.
While here, collapse the adjacent reasoning_budget_tokens override to a
single json_value() call; json_value already falls back to the default on
a missing/null/wrong-type key, so the explicit body.contains() guard was
redundant. No behavioral change.
Add a unit test in test-chat.cpp that exercises this path via
oaicompat_chat_params_parse with a Qwen3 template (which the autoparser
detects as a thinking-capable model) and asserts the returned
llama_params carries the per-request reasoning_budget_message rather than
the server default.
* cleanup
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
