* qwen4exp: MTP (NextN) self-speculative decoding support
Wires the checkpoint's 2.6B MTP head into the existing MTP speculative
framework for the qwen4exp architecture, following the deepseek4
conventions: block_count includes the NextN tail, the draft pass walks
[n_layer - nextn, n_layer) through the same layer-loop body, and the
target pass emits the pre-final-mixer wide stream as result_mtp_embd.
- hparams: read nextn_predict_layers; the tail is forced non-recurrent
(the full-attention-interval pattern would classify it as GDN) and
inherits the last main QSA layer's compress ratio
- loader: the tail layers and seven mtp.* entry/mixer tensors load only
when an MTP stage is requested (TENSOR_SKIP otherwise), so nextn-
bearing GGUFs serve normally on plain configs
- graph: MTP entry = per-reference residual_linear_shared fusion (joint
RMS over the flattened hc*n_embd stream, shared fc_hidden per branch,
fc_embedding added as a residual to every branch); exit reuses the
hc_head-shaped mixer with dedicated weights; delta-net construction
skips the draft pass (the draft context has no recurrent slots)
- llama: qwen4exp joins the MTP-capable arch allowlist and
llama_mtp_tail_uses_layer_cache; MTP feature width is the wide stream
(n_embd * hc_mult); the nextn pooling path gets its required
result_embd tensor
Measured on a 125B-A6B GGUF carrying the head (RTX 5090 + 24-core Zen5
hybrid, -ncmoe 38): draft acceptance 0.93-0.99, per-position 95%/pos on
code at n_max=4; draft passes cost ~2ms wall (131-144 nodes).
* cuda: dispatch the ABS and SGN unary ops
The kernels exist in unary.cu but neither the execution switch nor
supports_op routed them, so the qwen4exp PLE gate split to the CPU
scheduler twice per forward pass once its layer is offloaded (the cost
is the round trip, not the arithmetic — noted in the port's PR text).
* refactor(model): add qwen4exp NextN mixer tensors
* fix(qwen4exp): resolve standard NextN head layouts
* qwen4exp: decline per-step spec checkpoints (PLE conv tail not covered)
The per-step checkpoint path sizes its buffers from the ssm_* hparams
alone, but qwen4exp appends the PLE n-gram convolution history to the
tail of the same recurrent state row. A rejected draft token would roll
back the delta-net state and leave the PLE tail advanced. Decline
per-step (like openPangu) so the mode resolves to the whole-slot
gpu-fallback shadow, which covers the full row.
* qwen4exp: keep legacy mtp.* files loadable without MTP
Pre-rework files carried the NextN head under fork-local mtp.* names.
Register those seven names as known-and-skipped so such files still load
for non-MTP use instead of failing tensor accounting; an MTP request
against them reports the standard missing-layout error.
* loader: summarize skipped tensors in one line (per-tensor detail moves to debug)
An unused NextN/MTP head produces 33 ignore warnings per load; collapse
to a single summary at done_getting_tensors.
* spec: fall back to generic copy when checkpoint rows resolve no backend
Partial offload leaves host-resident recurrent state rows (CUDA_Host) matching
no backend's default buffer type between graphs, so checkpoint save/restore
asserted on a null backend. Drain the scheduler once and use the generic
tensor copy for unresolvable rows; the async fast path is unchanged. Also
guards the split-save branch that previously passed a null backend unchecked.
Diagnosed, patched, and validated by @joelfarthing (round-trip exact on
-ngl 6/8, -ncmoe control unchanged).
Co-authored-by: joelfarthing <joelfarthing@users.noreply.github.com>
* qwen4exp: support separate MTP head via -md (companion pair)
Lets a plain community quant load a predictor-only GGUF as the draft, so MTP
works without re-downloading a merged file and the head can sit on GPU at
full precision while the target runs partially offloaded. Relaxes the five
sites that assumed the target carries the NextN block: predict-layer count,
optional trunk tensors for mtp-only files, hc_attn_norm trunk marker,
n_embd_out derivation, and the wide-embd export predicate, plus qwen4exp
companion pairing enforcement.
Authored by @joelfarthing, replicated by @treo (3090, full 262K context).
Known limitation: llama-quantize on a predictor-only file trips the
attention-tensor count assertion; quantize before extraction.
Co-authored-by: joelfarthing <joelfarthing@users.noreply.github.com>
* server: name the stage chain in the single-slot refusal and say the fix
* qwen4exp: condense verbose NextN comments per CONTRIBUTING.md
* qwen4exp: accept the Dzannotti NextN head layout (predictor-only)
A predictor-only (-md) file may omit blk.N.nextn.shared_head_norm and carry the
head mixer norm only in output_hc_norm (the Dzannotti convention, with the +1
baked in as the direct-multiply grouped-RMS kernel expects). Fall back to
model.hc_head_norm as the head norm in that case, gated to mtp_only files where
output_hc_* is the head's own mixer and not the trunk's. Quimmedes-style raw
(non-baked) tensors remain intentionally unsupported.
Reported by @SamuelOliveirads.
* spec: merge the deepseek4/qwen4exp wide-stream feature-width branches
Both arches return n_embd_out under the same condition; one combined branch.
Suggested by @SamuelOliveirads.
---------
Co-authored-by: joelfarthing <joelfarthing@users.noreply.github.com>
Add LLM_ARCH_GEMMA4 to llama_model::supports_swa_compress() and route
gemma4's sliding-window mask through the compacted-cache builder, so
--swa-compress allocates the SWA layers at window size instead of at
full n_ctx.
Follows the per-arch pattern from #2253 (openpangu), #2266 (deepseek4)
and #2310 (laguna). gemma4's main graph path goes through
build_std_attention(), which #2310 already made compaction-aware, so the
only graph change is the mask and passing swa_head to the one direct
llm_build_kv() call on the shared-KV / no-wv branch.
The two hand-rolled builders in this file, build_gemma4_graph_parallel()
and build_gemma4_mtp(), are already unreachable with --swa-compress
(split/replicated caches and MTP are refused in llama_kv_cache_init and
llama_new_context_with_model respectively). Assertions are added to
document that rather than to change behaviour.
Measured on Goetia-26B-A4B (gemma4 26B-A4B, 30 layers: 25 sliding-window
with n_swa 1024, 5 global), CPU-only, -fa on -rtr -c 262144 -ub 2048:
flag off --swa-compress
KV self size 56320.00 MiB 5720.00 MiB
model + cache 74063 MiB 23463 MiB
prefill, 20530 tok 439.7 s 441.6 s (-0.4%)
decode, 200 tok 14.7 t/s 15.0 t/s (+2.4%)
Greedy output at temperature 0 is byte-identical with and without the
flag over a 20530-token prompt plus 200 generated tokens, which spans
about nine compaction cycles. With the flag off, pp512/tg32 are unchanged
from the unpatched build.
Co-authored-by: Alex <azilber@llme.consultent.ltd>
Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
Both the `-crs` help text and docs/parameters.md describe the parameter as a
"max of similarity ... that triggers prompt cache". It is the opposite: a
candidate cache entry is REJECTED when it falls below the value.
server_prompt_cache::load() takes it as `min_reusable_fraction` and applies it as
const float f_keep_cur = float(lcp_cur.first) / tokens.size();
if (f_keep_cur < min_reusable_fraction) {
continue;
}
so raising the value makes reuse stricter, while the current wording reads as
though raising it would admit more. The parameter name in the function signature
already says "min"; only the user-facing strings disagree.
Also states what the fraction is OF, which neither string did: the matched
prefix over the length of the CACHED entry, not over the new prompt.
Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
v_offset is used as a column index into the dequantized f16 buffer, but was
computed as a half-pointer difference. That is only correct for f16 K/V. With
q8_0 (34 bytes per 32 elements) it yields the wrong column, so the sparse
attention path reads V from the wrong positions and generation degenerates.
The per-32-block extra bits select the iq3nl_values half for each
16-element half: sub-block 4g+m, half h uses bit 8g+2m+h. The four
streams need local bits 0, 2, 4, 6, not 0, 1, 2, 3.
Previously --log-file only captured LOG()/LOG_TEE() macro output; the
LLAMA_LOG_* engine lines, server_log() output, and common_log (SLT_*/SRV_*)
slot/checkpoint lines all went to stderr only.
Route all three sinks to --log-file:
- llama_log_tee_callback tees raw llama/ggml output to the file and stderr
- server_log() mirrors its stdout line to the file
- common_log_set_file_ptr() shares LOG_TARGET's FILE* with the common_log
worker, which natively tees to stderr and file (so the SLT_*/SRV_* macros
stay untouched and bare LOG_*/QUE_*/RES_* calls are captured too)
Gate every sink on log_target_changed() (set only by log_set_target_impl(),
the wrapper --log-file calls) plus a != stdout && != stderr guard mirroring
LOG_TEE_IMPL. LOG_TARGET is non-null by default (log_handler() lazily opens
llama.log), so a plain null check would capture every line into a surprise
llama.log with no flag, and on a non-writable cwd would double-print to
stderr. log_target_changed() is marked at the wrapper rather than inside
log_handler1_impl because --log-file is parsed before any LOG() call, making
the first invocation's filename comparison vacuously false.
common_log_set_file_ptr shares the already-opened FILE* rather than calling
common_log_set_file, whose own fopen("w") would open a second handle on the
same path and let the two writes corrupt each other.
ggml_metal_init leaves ctx->encode_async nil, but
ggml_backend_metal_graph_compute invokes it unconditionally, so a Metal
backend created without a prior set_n_cb call segfaults on its first
graph. This affects rpc-server and ggml_backend_reg_metal_init; the
llama.cpp path is unaffected, since llama_graph_compute sets n_cb
before every compute.
Upstream added the same call in cad341d88 (#9698), the commit that
introduced encode_async. ik_llama.cpp carries the block but not the
initialization.
* rpc: disable unsafe memcmp graph cache
Same-shape prefill micro-batches compared equal and took the GRAPH_RECOMPUTE
path, re-running a stored graph against a grown KV context; the GLM-5.2 DSA
indexer then read past its buffers and crashed the server. Always send the full
graph. Upstream retired this cache design in ggml-org/llama.cpp#22701.
* rpc: use 64-bit ne/nb in rpc_tensor wire struct
ggml_tensor holds int64 ne and size_t nb; the wire struct stored them as uint32,
truncating any stride >= 4 GiB. The GLM-5.2 DSA indexer query stride crosses that
at ~26k tokens. Bump RPC_PROTO_MAJOR (wire-format change).
* ggml: use 64-bit locals in ggml_permute
Permuted strides were built in int locals, truncating any stride > 2 GiB before
it reached result->nb (size_t). Affects any permuted tensor over ~2 GiB.
* cuda : add the missing compile definitions to the HIP build
GGML_CUDA_FUSION, GGML_CUDA_MIN_BATCH_OFFLOAD and GGML_CUDA_PEER_MAX_BATCH_SIZE
are used unconditionally in common.cuh but were only defined in the CUDA branch,
so every HIP translation unit failed to compile. GGML_CUDA_IQK_FORCE_BF16 and
GGML_CUDA_F16 are user facing options the HIP branch silently ignored.
Also define GGML_USE_HIP. 43 tests in the sources imported from upstream use that
spelling, this fork only defined the older GGML_USE_HIPBLAS, so all of them took
the NVIDIA branch. Without it mmq_id_common.cuh defines TURING_MMA_AVAILABLE,
AMPERE_MMA_AVAILABLE, CP_ASYNC_AVAILABLE and FP16_MMA_AVAILABLE, i.e. the inline
PTX paths, and mmq_id.cu and mmq-instance-q6_k_id.cu then fail to build on
mma_new.cuh:181.
* cuda : shim the warp sync primitives in the HIP vendor header
ROCm 6 and later provide __shfl_sync() and friends as templates that static_assert
on the width of the mask, since an AMD wave can be 64 lanes wide. Define
HIP_DISABLE_WARP_SYNC_BUILTINS so that the shims replace them instead of clashing
with them, and add the shims the header was missing.
This is what makes mmq_id_common.cuh compile, so it unblocks all 26 mmq-instance-*_id
translation units, i.e. the MoE mat-mat path for the iqk quants.
* cuda : update the HIP vendor header for ROCm 6 and 7
- map nv_bfloat16 and nv_bfloat162 onto the __hip_bfloat16 types. Nothing declared
them, so every translation unit that mentions bf16 failed, convert.cu included -
that is the dequantize path the iqk quants use for prompt processing.
- use the hipblasComputeType_t and hipDataType entry points from ROCm 6.5 onwards.
hipblasDatatype_t is deprecated there and no longer matches hipblasGemmEx().
- make cudaStreamWaitEvent object-like. As a 3 argument function-like macro it did
not expand at the 2 argument call sites in reduce.cu, and the unexpanded name was
then passed on to CUDA_CHECK.
- add the mappings for cudaOccupancyMaxActiveBlocksPerMultiprocessor, which 29 of
the 30 failing flash attention translation units needed, and for the entry points
used by dsa_attn.cu and solve_tri.cu. hipBLAS spells a half as an unsigned short,
so cublasHgemmStridedBatched goes through a small casting wrapper rather than a
plain rename, which would not compile at the dsa_attn.cu call site.
* cuda : fix the remaining HIP build errors
- argsort.cu declared the sort order inside #ifdef GGML_CUDA_USE_CUB and used it
outside, and indexer_topk.cu calls argsort_f32_i32_cuda_cub() unconditionally
while it is only defined when CUB is available. Both break any build without CUB,
which includes MUSA and CUDA older than 11.7, not just HIP.
- ggml_backend_cuda_invalidate_graphs() touched ctx->cuda_graphs, which only exists
under USE_CUDA_GRAPH. The function is exported and called from llama-reload.cpp,
so guard the body rather than the function.
- solve_tri.cu included <cublas_v2.h> directly. common.cuh already pulls in whichever
vendor header is right, so the include is removed - it was redundant on CUDA too.
- dsa_attn.cu passed a data type where the GEMM entry point wants a compute type.
hipblasGemmStridedBatchedEx() has no data type taking overload. cuBLAS does, but it
is deprecated: cublas_api.h migrates CUDA_R_32F to exactly CUBLAS_COMPUTE_32F unless
the handle is in CUBLAS_PEDANTIC_MATH, which nothing in this tree sets. So this also
moves the CUDA build onto the primary entry point and drops a cublasGetMathMode()
per call. ggml-cuda.cu already passes a compute type to the same function.
- two mmvq instances called __dp4a() directly instead of ggml_cuda_dp4a(), the
wrapper the rest of the backend uses. On CUDA the wrapper is __dp4a() for every
architecture that has it.
- cap the flash attention vec f32 kernel at 4 columns per block on HIP. Both
logit_softcap variants of the 8 column kernel in one module overflow the 16 bit
branch offset of the AMDGPU backend.
* cuda : build the missing template instances in the HIP build
The HIP source list had drifted from the CUDA one and left out three families of
template instances that the backend references unconditionally:
- mmvq-instance-*.cu, the only definition site for the iqk mat-vec entry points.
iqk_mmvq.cu calls mul_mat_vec_iq4_ks_q8_1_cuda() and mul_mat_vec_iq4_kt_q8_1_cuda()
and nothing defined them, so the library did not link.
- the fattn-vec instances for q8_0-iq4_nl, iq4_nl-iq4_nl, q6_0-q5_0 and q8_0-q6_0,
which fattn-vec-f16.cu and fattn-vec-f32.cu dispatch to in the default
configuration.
- fattn-mma-*.cu. The MMA kernels are never selected on AMD, new_mma_available()
requires an NVIDIA device, but fattn-mma-f16.cu still needs the symbols.
* cuda : recognise AMD GPUs in GGML_CUDA_CC_IS_NVIDIA
CC_OFFSET_AMD is 1000000 and CC_OFFSET_MTHREADS is 0x100000, i.e. 1048576, so the
whole AMD range sits below the Moore Threads offset and every AMD GPU tested as
NVIDIA. turing_mma_available() then returned true on RDNA, the host picked an MMQ
tile of 128 while get_mmq_x_max_device() caps at 64 on AMD, and mul_mat_q_id hit
its NO_DEVICE_CODE guard and wrote NaNs. MUL_MAT_ID on IQ4_KS and IQ4_KT failed
this way on gfx1101.
No effect on CUDA, where a compute capability is 100*major + 10*minor and is
always far below CC_OFFSET_AMD.
* cuda : use v_perm_b32 for the 4 bit table lookup on HIP
HIP implements __byte_perm() in software: it stores an 8 byte union and does four
dynamically indexed byte loads, which end up in scratch. get_int_from_table_16()
calls it eight times per 32 weights, so every quant with a value table was paying
for that, while the trellis types were not.
__builtin_amdgcn_perm() is v_perm_b32, one instruction, and does the same job.
Taken from ggml-org/llama.cpp, which already carries this path.
Token generation on a 7800 XT, pure quantized Qwen2.5-1.5B, tg128:
IQ4_KS 16.30 -> 255.88 t/s
IQ4_XS 17.28 -> 268.05 t/s
IQ4_KT 201.76 -> 195.68 t/s (no table, unchanged)
Perplexity is unchanged to every printed digit and still matches the CPU exactly.
The function is duplicated in vecdotq.cuh and iqk_mmvq_templates.cuh, so both
copies need it - the iqk mat-vec instances only see the latter.
* cuda : use the shared flash attention support check on HIP
supports_op() carried a hand-rolled head size test for HIP that predates the
shared check: it accepted head size 64 with an f16 K cache and head size 128,
and nothing else. Head size 256 was rejected outright, so Gemma-2 and every
other 256 wide model fell back to the CPU for attention even though the
instances are compiled. @hardWorker254 reported 256 working with
ROCm 7.2.4 for both the f16 and the q8_0 cache.
Rather than adding 256 to the list, drop the branch and call
ggml_cuda_fattn_is_supported() as every other backend path does. It already
handles AMD: for cc >= CC_OFFSET_AMD it defers to the vec f16 or vec f32
support predicate depending on precision, which is exactly what
ggml_cuda_flash_attn_ext() dispatches to on AMD, because fast_fp16_available()
is true across the whole AMD cc range. The two now cannot drift.
This also removes a latent abort. The hand-rolled test returned true for any
head size 128 case regardless of the K and V types, so a combination without a
compiled instance, q4_1/q4_1 in a default build, reached the dispatcher and hit
on_no_fattn_vec_case() -> GGML_ABORT instead of falling back to the CPU. The
shared predicate is derived from the instances the build actually contains, and
after the source list repair earlier in this series the HIP build compiles the
same set as the CUDA build.
Beyond head size 256 this also lets HIP claim the asymmetric 192/128 and
576/512 vec f32 paths under GGML_PREC_F32. Those are untested on AMD; they are
gated by the same predicate CUDA uses.
* cuda : test the V head size, not the KV head count, for 192/128 vec f16 FA
ggml_cuda_fattn_vec_f16_is_supported() gates the asymmetric Dk != Dv branch on
if (K->ne[0] != 192 || V->ne[2] != 128) return false;
but ne[2] on K and V is the number of KV heads, not a head size. The test was
meant to be V->ne[0], as the wmma predicate added in the same commit (0459f595)
already writes it:
if (K->ne[0] != V->ne[0]) return K->ne[0] == 192 && V->ne[0] == 128;
and as the f32 twin has written it since 72201359 reworked that branch for
576/512. Only the f16 copy was left behind.
The kernels are there: ggml_cuda_flash_attn_ext_vec_f16() dispatches
FATTN_VEC_F16_CASE_DKDV(192, 128, ...) for f16-f16 and q8_0-q8_0 in both the
default and the GGML_CUDA_FA_ALL_QUANTS configuration, and the corresponding
hs192 instances are in the source list either way. The predicate just never
reported them, so a 192/128 shape whose KV head count was not coincidentally
128 was declined and attention fell back to the CPU.
This belongs in this series because the previous commit is what makes the
predicate reachable on AMD: with supports_op() routing flash attention through
ggml_cuda_fattn_is_supported(), the cc >= CC_OFFSET_AMD branch selects this
predicate for the default precision at every batch size, matching what
ggml_cuda_flash_attn_ext() dispatches to there. Without the fix the HIP build
would trade one hardcoded head size restriction for another.
NVIDIA is unaffected either way. Volta and later route 192/128 through the mma
or wmma predicates, and on Pascal the Q->ne[1] <= 8 decode case is diverted to
vec f32 before this predicate is consulted.
The DSA attention kernel used the shared cublas handle without binding
it to the backend's stream, so its Q.K / P.V GEMMs ran on a different
stream than the gather and softmax kernels. The softmax could then read
the score buffer before the GEMM wrote it, picking up stale (NaN) values.
* vulkan : use ggml_row_size for types with a per-row scale
Types that declare a row_meta_size store a per-row scale ahead of the row's
blocks, so a row is not ggml_type_size()*ne/ggml_blck_size() bytes. This
under-sized src0 in the four quantized mat-mul paths, and made
ggml_vk_dim01_contiguous() report such a tensor non-contiguous, which in turn
made supports_op reject it. No change for row_meta_size == 0.
* vulkan : add IQ4_KS and IQ4_KT support
A row of these types is one f32 scale followed by the row's blocks, so rows are
not a whole number of blocks apart and the usual block-indexed addressing does
not work. They are read through a uint32_t alias of binding 0 and addressed by
word; types.comp holds the alias, the stride and the decode, so each shader only
expresses its own addressing and no push constant layouts change.
Covers to_fp16, get_rows, mul_mat_vec (incl. MUL_MAT_ID) and scalar + coopmat1
mul_mm. get_rows addresses by row rather than through nb01/02/03, which cannot
express a per-row scale, so supports_op accepts only a contiguous src0 for these
two types. coopmat2 is excluded because coopMatLoadTensorNV addresses through a
uniform grid tensor layout, which cannot describe the row prefix; the two
mat-mat getters return nullptr there and the callers fall back to F16.
* tests : add IQ4_KS/IQ4_KT decode validation
Re-implements in C++ the indexing each of the four shader families uses and
diffs it against ggml's to_float over several row and block counts. CPU only: it
validates the format transcription, not the compiled shaders.
* Load standalone Qwen3.5 MTP GGUFs passed with -md
A predictor-only MTP GGUF reports the full block count (n_main +
nextn_predict_layers) but only ships the NextN block, so loading one
with -md failed:
check_tensor_dims: tensor 'blk.0.attn_norm.weight' not found
create_qwen35_tensors() and create_qwen35moe_tensors() create every
main block as required. Detect the predictor-only case the same way
create_step35_tensors() does and mark the absent blocks
TENSOR_SKIP|TENSOR_NOT_REQUIRED.
Qwen3.5 also has to use the common MTP package contract, otherwise the
predictor-only GGUF is never classified as a companion, and the target
is not classified TARGET_ONLY - which is what makes it export the
hidden states the companion consumes.
The remaining two hunks cover cases the above newly reaches: a
predictor-only GGUF passed as -m now loads far enough to abort in the
graph builder, and its empty main blocks reach split_recurrent_tensors()
under -sm graph.
* Qwen3.5 MTP: require q_proj in predictor-only GGUFs, check companion arch
Review follow-up.
A dense NextN block loads q_proj as optional because it can be shared
with the last main block. A predictor-only GGUF has no main blocks, so
one built that way loaded with wq == nullptr and then hung. Require the
tensor in that case so the load fails naming it. eh_proj, attn_q and the
MLP are all optional on that block, so the tail probe stays on enorm,
which is required - the comment there said only eh_proj.
Adding Qwen3.5 to the common MTP package contract also made
common_speculative_has_recognized_mtp_companion() accept any GGUF
classified COMPANION, with no architecture check of the kind the Step
and DeepSeek branches have. Add it, plus the predictor count. Dense and
MoE are separate architectures, so the comparison is on the arch itself.
OpenAI-compatible clients (e.g. pi coding agent) send
max_completion_tokens for the output token cap on custom
openai-completions providers. The server only read n_predict and
max_tokens, so the cap was silently dropped and n_predict fell back
to -1 (unlimited). This allowed runaway generations of 40k+ tokens
on long agent sessions.
Matches upstream llama.cpp behavior where max_completion_tokens is an
alias of n_predict (tools/server/server-schema.cpp).
Marvin
Joel Farthing
Samuel Oliveira Alves
James
Alexey Zilber
Kawrakow
mattafaak
Alex
Yap Sok Ann
Skelectric
Riccardo Chiumiento
Coenie Beyers
Guy Barel
dungquixote42
Thireus ☠
Jun Yamog