* 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>
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.
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.
* model: Ling-3.0 (bailingmoe3) runtime support
* model: Ling-3.0-tiny support
---------
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
* Adding Muse-Glimmer support
* Different rms_eps for post norm ops
* Need attn_post_norm split for Muse-Flimmer
* WIP: split mode graph
* Forgot this file
* Clean it up
* Minor
* Muse-glimmer: Slightly better split mode graph (+2% TG)
* CUDA indexer topk: this is better for PP
* Don't overstep
* Cleanup
* Allow Q8_0 cache in the CUDA DSA implementation
* DS4: do not cast caches to f32
* Fix massive inefficiency in CUDA Q->f32/f16 and f32/f16->Q copies
* Re-enable -ictk | --indexer-cache-type-k
* ggml : assert the delta-net value head dim equals the key head dim
The CPU forward sizes its result from the value head dim (src[2]->ne[0]) but
indexes it with the key head dim (src[0]->ne[0]). A model where the two differ
was mis-indexed silently. The CUDA op has asserted this all along.
* ggml : fuse the delta-net recurrent state copy into the op
The problem: the delta-net op produces the new recurrent state into the tail of
its result. Then llama copies that tail into the KV slot the state was read
from. The copy buys nothing - the kernel could write the slot itself.
The change: the slot's two halves are written by two narrow CPY nodes instead of
one CONCAT. That isolates the state write in an ordinary node, which the
scheduler places by the rules it already has. A backend that recognises the
pattern lets the kernel write the slot directly and skips that node. One that
does not implement the fusion runs the copy as before. No public header changes.
Notes: this leaves ggml_concat_inplace(), added in #1777 for exactly this site,
without a caller.
Marvin
James
Alex
Kawrakow
Yap Sok Ann
Riccardo Chiumiento
Coenie Beyers
Joel Farthing
Guy Barel
Petr Vilím
Nexesenex