TEMP integration: compose #131 qwen35moe with current ROCm shadow - #35
Merged
nekomario28 merged 34 commits intoAug 26, 2026
Merged
Conversation
The vendored sgl-kernel/llama.cpp CUDA kernels under csrc/gguf/ already
dispatch 19 quant types, but the Python side only ever described 6, so
K-quants and I-quants were unreachable. This wires the Python tables up to
what the kernels actually implement.
- csrc/gguf/gguf_kernel.cu: add default: TORCH_CHECK guards to all five
switch (type) blocks. Y is allocated with torch::empty, so an unsupported
type previously returned uninitialized garbage rather than raising;
ggml_moe_get_block_size returned 0. Also null-check the ggml_get_to_cuda
function pointer, which ggml_dequantize called unconditionally.
- models/gguf/dequant.py: extend GGML_* constants, BLOCK_SHAPE and GGML_NAME
to all 22 types (block geometry derived from ggml-common.h with QK_K=256,
K_SCALE_SIZE=12). Add five frozensets mirroring the C dispatch switches,
each citing the file:line it mirrors. The pure-torch reference dequantize()
still implements Q4_0/Q6_K only and now says so explicitly.
- layers/gguf.py: replace the three hardcoded {Q4_0, Q8_0, Q6_K} sets with
the shared tables. Because _MMQ was tested before _DEQUANT and the sets
were identical, the dequant branch was unreachable dead code; it is now
the live path for I-quants, which have MMVQ and dequant kernels but no
MMQ kernel. This is what makes I-quant prefill work at all.
- moe/fused_q4_0.py: add fused_experts_gguf(quant_type), with
fused_experts_gguf_q4_0 kept as a behaviour-preserving wrapper.
- moe/cpu_executor.py: raise a clear NotImplementedError for GGUF formats
other than q4_0. WF_Q4_0 is the only GGUF format with an AVX/VNNI kernel
in cpu_moe_ext.cpp; the limit was previously implicit.
Tests cover the tables against ggml-common.h and the capability sets against
the C switches themselves, so the two cannot drift.
Not covered: per-architecture GGUF shims (GGUF_ARCH_TO_REGISTRY still lists
gemma4 only) and CPU AVX kernels for non-Q4_0 experts. Untested on hardware.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
Builds the FreeToken ModelConfig from llama.cpp qwen35moe GGUF metadata and maps tensor names back to FreeToken parameter names. Derived against vcruz305/Ornith-1.5-35B-A3B-GGUF. The GDN projections are the non-obvious part. llama.cpp's qwen3.5 mapping splits what qwen3next fused -- attn_qkv<-in_proj_qkv, attn_gate<-in_proj_z, ssm_beta<-in_proj_b, ssm_alpha<-in_proj_a -- and FreeToken's HF loader already fuses those same pairs (_PT_FP8_FUSE / _PT_BF16_FUSE in weight.py), so we emit the fused representation and the model code sees one shape regardless of source. Head layout is pinned by arithmetic rather than assumed: attn_qkv's packed output width of 8192 equals q(group_count*state_size) + k(group_count*state_size) + v(time_step_rank*state_size) = 2048+2048+4096, so num_k_heads==ssm.group_count and num_v_heads==ssm.time_step_rank at head_dim==ssm.state_size==128. parse_gguf_config asserts the inner_size relation rather than trusting it. Deliberately NOT added to GGUF_ARCH_TO_REGISTRY yet: iter_gguf_weights is not implemented, and the routed-expert path cannot serve this checkpoint regardless -- Ornith's experts are IQ3_S (gate/up) + Q4_K (down) as separate stacks, while the offload/CPU bank machinery keys on a single fused "q4_0" schema. Registering the arch now would advertise support that does not exist. See the follow-up plan. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
…ommon.h The header-derived size check dropped block_iq3_s's scales[IQ3S_N_SCALE] field because the macro table lacked that define, yielding 106 bytes and failing against BLOCK_SHAPE's 110. 110 is correct and independently confirmed against a real checkpoint: Ornith-1.5-35B IQ3_M's token_embd.weight is 248320*2048/256*110 = 218,521,600 bytes on disk, matching row_bytes() exactly. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
Adds the architecture path plus the mixed-quant machinery llama.cpp checkpoints actually need. GGUFMergedLinear (layers/gguf.py): a merged projection whose parts carry different ggml types. FreeToken fuses projections by concatenating packed rows, which needs a shared row_bytes -- true for gemma4 (all Q4_0) but not for llama.cpp's mixed quants. Ornith's full-attention qkv is q,k IQ3_S + v Q4_K and its GDN in_proj is attn_qkv Q4_K + gate/beta/alpha IQ3_S; 880- and 1152-byte rows cannot interleave. Concatenating the per-part OUTPUTS is identical to the merged GEMM since every part reads the same x. gguf_merged_or_plain() keeps the uniform case on the cheaper single-kernel GGUFLinear. qwen3_5_moe/gguf.py: iter_gguf_weights + convert_qwen35_to_gguf. Maps the GGUF tensor table onto the module tree, fusing per layer in the orders gdn.py and attention.py declare (in_proj = attn_qkv, attn_gate, ssm_beta, ssm_alpha, matching _in_proj_split = [conv_dim, value_dim, n_v, n_v]). A_log/dt_bias stay fp32 as gdn.py requires; conv1d is reshaped to [conv_dim, 1, kernel] for its .squeeze(1). qwen3_5_moe/gguf_experts.py: routed-expert host banks in native packed bytes. Offload path generalized off the q4_0 literal: a "gguf" bank schema, byte accounting from row_bytes, and the (gate_up, down) ggml types threaded ModelConfig -> ExpertBanks -> OffloadMoeCache -> the MoE kernels, because a GGUF's row stride is a property of the file, not the format. fused_experts_gguf takes a separate down_quant_type: the two banks own separate slot pools and may legitimately differ. What is NOT supported, and fails loudly rather than silently: a bank whose ggml type varies BY LAYER. moe_vec.cuh addresses the pool as expert * nrows * (ncols / qk) with no padding allowance, and the pool is one allocation shared by all layers, so mixed strides would read every block at the wrong offset and emit fluent garbage. _gguf_banks raises with the offending layers named. Ornith IQ3_S and IQ3_XXS have uniform banks; IQ3_M, IQ2_M, IQ2_XXS and IQ1_S split ffn_down_exps across two types. Also: _model_setup_override now declines for expert_quant == "gguf". qwen3_5_moe exports an NVFP4 setup_offload_expert_banks that would otherwise hijack the GGUF load and fail deep inside an NVFP4 reader. Not yet run end-to-end; weights have not been loaded on hardware. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
_LAYER_MAP and _FUSE were written before the module tree was known and mapped attn_q/k/v -> self_attn.q_proj/k_proj/v_proj and ffn_gate_shexp/ffn_up_shexp -> mlp.shared_expert.gate_proj/up_proj. Those attributes do not exist: Qwen3_5Attention builds one merged qkv_proj (_qkv_split [8192, 512, 512]) and _SharedExpert one merged gate_up_proj. iter_gguf_weights always handled them by fusion, so the entries were dead -- but gguf_name_to_freetoken would hand a caller an invented parameter name. Replaced both with _MERGED_PARTS, the set of suffixes that are parts of a merged projection, which gguf_name_to_freetoken now reports as None (handled elsewhere). The constant documents each merged target and its concat order, including why the GDN target is in_proj rather than in_proj_qkvz/in_proj_ba (gdn.py only splits those out on the fp8 branch, and a GGUF checkpoint is not fp8). Also: tests call .forward() explicitly -- BaseOP defines no __call__. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
…names The name-map test asserted attn_q/k/v -> self_attn.q_proj/k_proj/v_proj and ffn_gate_shexp/ffn_up_shexp -> shared_expert.gate_proj/up_proj. Those attributes do not exist on the module (one merged qkv_proj and one merged gate_up_proj), so the test was locking in names the loader could never populate. Assert None for fusion parts instead, which is the contract iter_gguf_weights actually implements. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
…type+shape Three defects the module-tree reconciliation caught on the real checkpoint: 1. The shared-expert suffixes (ffn_gate_shexp / ffn_up_shexp / ffn_down_shexp) were handled AFTER the per-layer-kind branch, whose `else: continue` swallows any suffix it does not recognize. Every layer's shared_expert.gate_up_proj and down_proj therefore went unfilled -- 81 buffers -- while the loader reported success. Moved the handling above that branch, where both layer kinds reach it. 2. mlp.shared_expert_gate: llama.cpp stores the single-output gate as a 1-D [hidden] vector; LinearReplicated(hidden, 1) declares [1, hidden]. Now reshaped. 3. linear_attn.conv1d.weight was yielded fp32. gdn.py exempts only A_log and dt_bias from the model-dtype downcast, so the depthwise conv is bf16. Found by reconciling every name/shape/dtype iter_gguf_weights yields against the constructed module's own state_dict, which is the check that catches this class of bug before it becomes fluent garbage at generation time. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
…. filter)
iter_gguf_weights handled token_embd and output_norm as globals, then bailed on
`not name.startswith("blk.")` -- which silently discarded output.weight, the untied LM
head. lm_head.qweight [248320, 1680] stayed at its torch.empty allocation, so the model
would have generated from uninitialized memory. Ornith ships output.weight (hence
tie_word_embeddings=False), so it is a real buffer that must be filled.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
…c_layer_banks
Both bugs surfaced only on IQ3_S -- IQ3_M happened to mask them.
convert_qwen35_to_gguf hardcoded quant types for o_proj (Q4_K), out_proj (IQ3_S), the
token embedding and the lm_head, and used quant_map.get(key, DEFAULT) elsewhere. Q4_K is
right for IQ3_M's attn_output but IQ3_S uses IQ3_S there, so every full-attention o_proj
was allocated 2304 bytes/row instead of 1760. A guessed quant type silently mis-sizes a
packed buffer and the only symptom is garbage output, so the new qt() helper raises on a
missing tensor instead of defaulting. _scan_quant_types now also keys the globals
(token_embd/output) under layer -1 so the embedding and head are sized from the file too.
gguf_expert_specs returned {name: [per-layer (shape, dtype)]}, but alloc_layer_banks
takes {name: (shape, dtype)} -- one spec per bank -- so bank allocation died with
"too many values to unpack". The per-layer form was wrong by construction anyway: a
bank's slot pool is one allocation with one stride, so its type must be uniform. It now
returns the flat form and rejects a non-uniform bank.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
…ansformers)
transformers keys its GGUF tokenizer converters by its own model_type, so passing the
GGUF arch string straight through raised KeyError('qwen35moe') and killed the detokenizer
worker during load. qwen35moe is a GPT2-style BPE with merges (tokenizer.ggml.model=gpt2,
pre=qwen35), which the existing qwen2 converter handles.
Two adjacent gemma4-isms were also wrong for any other vocab and are now per-arch:
- eos preferred the literal <turn|>, and gguf_eos_token_ids hardcoded <eos>/<turn|>.
Neither exists in Ornith's vocab, so generation would have used the wrong eos and
registered no chat stop id -- every request running to max_tokens. _STOP_TOKENS maps
arch -> stop names in preference order; gemma4 keeps (<turn|>, <eos>) so its behaviour
is unchanged, qwen35moe gets (<|im_end|>, <|endoftext|>), matching the file's
eos_token_id 248046.
- unk defaulted to the literal <unk>. Ornith has no unknown_token_id and no <unk> token,
and handing PreTrainedTokenizerFast a token absent from the vocab appends it. tok_for
now drops a default that is not in the vocab.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
…ed_tokens.weight
convert_qwen35_to_gguf existed but nothing called it, so the module tree still held dense
ops and load_state_dict raised KeyError('model.embed_tokens.weight') against a loader that
yields .qweight. gemma4 performs its swap at the end of the model's __init__ behind
is_gguf_model(); qwen3_5_moe now does the same.
Unlike gemma4 -- whose checkpoint is uniformly Q4_0/Q6_K, so its swap needs no file access
-- this swap sizes every packed buffer by the ggml type that tensor actually uses, which
only the file knows. parse_gguf_config therefore publishes ModelConfig.gguf_model_path and
__init__ asserts it is set.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
Threading the bank types through the engine, I bulk-replaced both `quant_format=banks.quant_format` sites -- but one of them is resolve_moe_cache_auto(), which only sizes the slot pool from bank bytes and has no such parameter. Only the OffloadMoeCache constructor takes it. TypeError at cache-sizing time, after the weights and expert banks had already loaded. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
…t garbage The MoE expert banks were built as cat([gate, up], dim=0).reshape(E, 2I, rb). gate and up each arrive from the reader as [rows, row_bytes] = [E*I, row_bytes(H)], and ggml's fastest-first dims [H, I, E] make E the slowest axis, so those rows are EXPERT-MAJOR. Concatenating on dim 0 therefore lays down every expert's gate before any up: expert 0's slice became its own gate rows plus expert 1's gate rows, and its up rows sat E*I rows away. Every routed expert in all 40 layers read another expert's weights. This loaded cleanly, ran at full speed (56 tok/s decode on a 4060), and generated confident nonsense -- 'r med media media media middle med med middle med...' -- which is why the fix comes with a byte-exact identity check rather than a reasoned argument: build the bank, then assert bank[e, :I] is gate rows [e*I, (e+1)*I) and bank[e, I:] is that same expert's up rows, for experts at both ends and the middle. The old code passes for expert 0's gate and expert 255's up and fails everywhere between, which is exactly the interleaving signature. Fixed by reshaping each part to [E, I, rb] and concatenating on the row axis (dim=1), so each expert's gate and up stay together. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
Qwen3.5 stores these norm weights as w with an effective scale of 1+w, and FreeToken's GemmaRMSNorm multiplies by the raw weight -- so weight.py bakes the +1 in at load time (_is_gemma_norm / _GEMMA_NORM_SUFFIXES). The GGUF loader yielded them raw, so every RMSNorm in all 40 layers scaled by w instead of 1+w. w is typically 0.01-0.1, i.e. the residual stream was attenuated 10-100x at every norm, which is why generation was fluent nonsense rather than an error. Shifted: model.norm, input_layernorm, post_attention_layernorm, self_attn.q_norm, self_attn.k_norm. NOT shifted: linear_attn.norm (ssm_norm) -- the GDN gated norm is a standard w*x norm and weight.py excludes it too. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
Two value-level bugs, both invisible to structural checks -- shapes, dtypes, byte identity and per-layer activation magnitudes were all correct while the model emitted fluent nonsense. 1. ssm_a. llama.cpp writes this tensor pre-transformed: it holds A = -exp(A_log), not A_log. Every value in Ornith IQ3_S is negative, spanning [-70.11, -0.0189]. gdn.py computes g = -A_log.exp() * softplus(a + dt_bias), so feeding it A gives -exp(-10.6) ~ -2.5e-5 and the recurrent decay gate collapses to zero in all 30 GDN layers. Now yields log(-A), so -exp(A_log) reproduces the stored A exactly, and raises if the tensor is not wholly negative rather than silently producing NaN. 2. The (1+w) RMSNorm shift added in the previous commit was wrong for GGUF and is reverted. HF stores raw w and weight.py adds 1.0, but llama.cpp's converter already folds the +1 into what it writes, so the shift double-counted. Measured means on Ornith IQ3_S: attn_norm 0.920, post_attention_norm 1.135, attn_q_norm 1.326, output_norm 2.640 -- all centred near 1, whereas raw w centres near 0.0X. Both were found by reading the actual stored values against what gdn.py and GemmaRMSNorm expect. Structural verification cannot catch a tensor that is the right shape and the wrong quantity. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
…cts HF order When num_k_heads != num_v_heads, llama.cpp's converter reorders every tensor that indexes the V-head dimension so ggml_repeat can replace an interleaved repeat (conversion/qwen.py, _LinearAttentionVReorderBase): HF / FreeToken (grouped by K head): [G0_v0, G0_v1, G1_v0, G1_v1, ...] GGUF (tiled for ggml): [G0_v0, G1_v0, ..., G0_v1, G1_v1, ...] FreeToken's GDN pairs K head k with V heads [k*R, (k+1)*R) -- grouped -- and the safetensors loader applies no reorder because HF is already grouped. So every K/V pairing in all 30 GDN layers was wrong. Ornith: 16 K heads, 32 V heads, R=2, head_v_dim=128. Inverted on all seven tensor groups llama.cpp touches: attn_qkv (V rows only), attn_gate, ssm_beta, ssm_alpha, ssm_a, ssm_dt.bias, ssm_conv1d (V channels only) and ssm_out (columns). Row permutations are applied directly to the packed bytes -- each output row is an independent run of quant blocks, so reordering rows never splits one. ssm_out is the exception: its COLUMNS carry the V dimension, and a 128-wide head straddles the 256-element IQ3_S blocks, so it cannot be permuted while packed. That one tensor is dequantized to dense bf16 (16 MiB per GDN layer, ~503 MiB total) and convert_qwen35_to_gguf leaves linear_attn.out_proj as a dense Linear. Like the A_log transform, this is invisible to every structural check: shapes, dtypes, byte identity against the source, and per-layer activation magnitudes are all correct while the K/V pairing is scrambled. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
…0/Q6_K only) _to_bf16 goes through models/gguf/dequant.dequantize, which implements only Q4_0 and Q6_K -- it is the reference/test path. ssm_out is IQ3_S, so materializing it dense for the column un-tiling raised NotImplementedError. _dequant_any routes through ggml_dequantize, which covers all 19 types, and returns a CPU tensor so the normal load path places it. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
…round The refusal said which layers disagreed but not why it is impossible or what to do instead. It now states that the slot pool is one allocation and moe_vec.cuh indexes it with no padding allowance, that llama.cpp's *_M and *_XXS levels raise the precision of the first few layers' ffn_down_exps, and that llama-quantize --pure produces a uniform checkpoint that loads. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
Same hybrid GDN/full-attention decoder as qwen35moe with a plain SwiGLU MLP instead of routed experts, so it reuses the package, the model classes and this adapter. - _kv() prefixes metadata keys with the checkpoint's own general.architecture rather than a hardcoded "qwen35moe", since the dense arch string is "qwen35". - parse_gguf_config reads feed_forward_length into intermediate_size when expert_count is absent, and leaves expert_quant "none" so the engine does not go looking for offload banks that do not exist. - is_gguf_model now keys on gguf_model_path instead of expert_quant, otherwise the dense path would skip the op swap entirely. - iter_gguf_weights maps ffn_gate/ffn_up/ffn_down onto mlp.gate_up_proj and mlp.down_proj (Qwen3_5DenseMLP subclasses _SharedExpert, so the same fused shape applies). Handled before the per-layer-kind branch, whose else: continue would otherwise drop them. - convert_qwen35_to_gguf derives the qkv and in_proj split widths from the config instead of Ornith's constants. Ornith-1.5 is 16 q heads / 2 kv / 32 v; Qwen3.8-27B is 24 / 4 / 48, so the hardcoded [8192, 512, 512] and [8192, 4096, 32, 32] were wrong for anything else. Dense checkpoints have no expert banks, so the uniform-bank restriction does not apply to them: Qwen3.8-27B Q4_K_M mixes Q4_K and Q6_K across layers for attn_qkv, attn_v and ffn_down, and per-tensor quant types already handle that. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
models.md still said native GGUF was Gemma-4 only, which stopped being true once qwen35moe and qwen35 landed. Adds a GGUF section listing the architectures by their general.architecture string, the quant types with the prefill and decode path per family, and the two constraints a user hits when picking a file: routed-expert banks must use one ggml type across all layers (llama-quantize --pure avoids it, dense models are exempt), and GGUF paths are TP=1 with any NextN/MTP block dropped.
iter_gguf_weights asserted include_moe_experts was False unconditionally, because the MoE variant serves its routed experts from the offload cache. A dense qwen35 checkpoint has no routed experts, so the engine correctly asks for everything and the assert tripped during load. Now conditional on the checkpoint actually having experts.
Same failure as qwen35moe had: convert_gguf_tokenizer is keyed on transformers' own model_type, so an unmapped GGUF arch string raises KeyError and kills the detokenizer worker during load. The dense variant shares the MoE one's tokenizer exactly (gpt2 BPE, pre=qwen35, 248320 tokens, eos <|im_end|>), so it maps to the same qwen2 converter and the same stop tokens. (cherry picked from commit e43bf64)
Multi-shard was the availability blocker: every large model ships split (-00001-of-000NN), so quant-type support was unreachable for exactly the checkpoints it mattered for. reader.py now resolves any shard (or a directory of shards) to the ordered set, reads metadata/arch/tokenizer from shard 1 only, and aggregates tensor tables across shards. Validated against llama.cpp's convention: split.no is 0-based while filenames are 1-based, and split.tensors.count is the TOTAL across shards, not per shard. A missing or extra shard raises with the indices named rather than loading a partial model. qwen3moe support reuses the qwen35moe adapter's shape with the GDN complexity removed: plain attention on every layer, no shared expert, no MTP, and no (1+w) norm shift (qwen3_moe uses RMSNormFused, and plain Qwen3 GGUFs are not pre-shifted). Merged qkv still routes through GGUFMergedLinear because Qwen3 quants mix types across q/k/v. Also fixes the same defect qwen35moe had: Qwen3MoeForCausalLM never called convert_qwen3moe_to_gguf, so the module kept dense ops while the loader yielded .qweight. Caught by the review pass, not by construction. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014h3QjryXQL6PXJWQdA6tvu
The sibling-shard glob was f"{base}-?????.of-{total}.gguf" with a literal dot where the
-of- separator belongs. ????? consumes the index, then the dot has to match the hyphen of
-of-, so the pattern matched nothing and a complete set reported every index missing.
Verified against a real 3-shard split of Ornith IQ3_S produced by llama-gguf-split.
Also resolves a directory of shards to its first shard, since a user is more likely to
point at the download folder than at a specific shard, and raises if the folder holds more
than one split model.
The generated tests called gguf.GGUFWriter(path) without its required arch positional, so six of seven failed on a TypeError before testing anything. Hand-writing the format keeps these independent of that writer's API and is only about forty lines. Covers the three properties of llama.cpp's split convention that are easy to get wrong and invisible when you do: split.no is 0-based while filenames are 1-based, split.tensors.count is the total across shards rather than per shard, and only shard 1 carries the metadata. There is also a fixture guard asserting shard 2 genuinely lacks general.architecture, since otherwise the shard-1-resolution test would pass vacuously.
Adds the qwen3moe row and the split-checkpoint paragraph to models.md, and updates the fork README's verified table with the sharded Ornith run and Qwen3-30B-A3B. Corrects the Ornith IQ3_XXS row to 6/6: the earlier 5/6 was a truncated preview in my harness, not a wrong answer.
Q4_0 was the only GGUF type with a CPU dot kernel, so --moe-backend cpu and hybrid refused every K-quant checkpoint and pushed those users onto offload. That is backwards: the people who need the hybrid backend most are the ones short of VRAM, and they are running exactly these quants. Adds scalar W4A16 dots for Q4_K and Q6_K, mirroring dequantize_block_q4_K and dequantize_block_q6_K in gguf/dequantize.cuh element for element. Correctness first; an AVX2/VNNI version is a separate job. Both skip the int8 activation path Q4_0 uses, since the super-block scale structure does not map onto it. Also wires quant_format == "gguf" through to a concrete format. The GGUF offload path tags the cache "gguf" and carries the ggml types separately, so the executor was matching on format strings that nothing produces and would have refused a real checkpoint even for Q4_0. _resolve_gguf_format maps the bank types onto one weight format and refuses with a specific reason when it cannot: mixed banks (Q4_K_M puts ffn_down_exps in Q6_K) get a different message from types with no kernel at all, because the fix differs. Tests check both kernels against the vendored CUDA dequant plus the production bf16 GPU decode on byte-identical banks, at three batch sizes. Blocks are synthesized rather than carved from a checkpoint so this runs without a multi-GB download. Worth recording what these tests caught, since neither bug is visible structurally and both give you a model that loads and runs at full speed: Q4_K quants are unsigned 0..15 with a per-sub-block min, not Q4_0's q-8, and a byte's high nibble is element l+32 rather than l+1. Q6_K's qh shift is per group, and an extra term in the lane index corrupted the upper half of every group while still correlating 0.17, which reads like a tolerance problem rather than a wrong kernel.
…tion A GGUF checkpoint with untied embeddings had its lm_head swapped for a plain GGUFLinear, which has no reason to know it is the head and so projected the whole prefill hidden state. ParallelLMHead and gemma4's GGUFTiedLMHead both already drop everything but the last position of each sequence, so only the untied GGUF path was affected, and only since the qwen35moe/qwen3moe adapters landed. Logits are [tokens, vocab] and on a large-vocabulary model that is the biggest tensor in the model. Ornith-1.5's vocab is 248,320, which in bf16 is 486 KiB of logits per token, so a 1,800-token prompt asked for a single 894 MB allocation. Every row but one was then thrown away by the sampler. Measured on an RTX 4060 Laptop (8 GB), Ornith-1.5-35B-A3B IQ3_S, offload backend. Before, a ~766-token prompt prefilled in 3.6s and ~1,800 tokens either hung past 300s or killed the worker with "CUDA driver error: device not ready". After: 1,834 tokens in 6.1s, 2,487 in 9.6s, 3,765 in 13.9s, 4,598 in 16.9s. Decode was never affected either way, which is what made this look like a memory problem rather than a head problem: generating 30,000 tokens works fine because decode only ever has one position per sequence to project. Adds GGUFLMHead rather than putting the slice in GGUFLinear, since that class is also every q/k/v/o and MLP projection in the model and must not slice. Beyond this fix, an 8 GB card still runs out of room around 5-6k prompt tokens, but in a different place: the GDN chunked kernel's activations (chunk_gated_delta_rule_fwd -> chunk_fwd_o). That is headroom rather than a defect and --max-prefill-length handles it; at 2048 the same box prefills 6,160 tokens in 23.2s and 7,784 in 27.5s. Worth knowing that the 8192 default does not chunk a typical long prompt at all.
Two problems, both only reachable on pre-Ampere cards. The head_dim <= 128 branch of _select_extend_tile returned (128, 64) unconditionally, with no fits() check at all, unlike every other branch. At block_d 128 that tile needs 96KB, which is fine on sm_80/sm_89 and fatal on Turing's 64KB hard limit. The shared-memory estimate itself also has the wrong shape for Turing. Ampere+ tiles with mma.sync m16n8k16, Turing with m16n8k8, and Triton stages materially more for the same logical tile as a result. Measured from triton's own OutOfResources reports on a Quadro RTX 6000 (65536 limit): (128,64)@block_d 256 -> 196608 (64,64)@256 -> 131072 (64,32)@256 -> 98304 (128,64)@128 -> 98304 The existing (block_m + 2*block_n)*block_d*2 estimate scales with N where the real cost scales with M, so it under-predicts exactly the tall-M tiles that overflow. A head_dim-256 model on a 64KB card therefore selected (64,32), which the estimate said fit at 65536, and still died needing 98304. The corrected bound is applied ONLY when compute capability < 8. The same tiles measure 114688 for (128,64)@256 on sm_89, so applying it there would shrink tiles that genuinely fit and cost roughly 2x on large head_dim. Ampere+ selection is unchanged, and the existing parametrised tests over A100/H100/ sm_89/unknown budgets all still pass. Turing also needs rungs below (64,32) at head_dim <= 256, so the ladder continues to (32,32) and (32,16) there; Ampere+ keeps its original two-way choice. Verified on the hardware: Ornith-1.5-35B-A3B IQ3_S on a Quadro RTX 6000 (sm_75, 24GB) now serves, where before it died at warmup. 64-69 tok/s with 10240 of 10496 expert slots resident. Correctness spot-checked at temperature 0 against known facts ("The capital city of France is" -> " Paris, ...").
moe_vec.cuh carries the routed-pair index (token * top_k + slot) in gridDim.z, which CUDA caps at 65535 on every architecture. All 19 launchers computed the grid as (block_num_y, 1, tokens * top_k), so a prefill wider than 65535 / top_k tokens failed the launch with cudaErrorInvalidConfiguration, which torch reports as "CUDA error: invalid argument". At top_k 8 the ceiling is 8191 tokens. fused_experts_gguf issues two launches, gate_up as (top_k, N) and down as (1, N * top_k), and both reach z = N * top_k, so both overflow together. Nothing about this is architecture- or OS-specific; it needs only a long enough prompt and was reported with a 20k-token one. Fixed by launching token-aligned chunks rather than reshaping the grid. Moving the pair index into x would also fit, since gridDim.x allows 2^31-1, but it would cost the locality the kernel is built around: with the row index in x, consecutive blocks walk rows of the SAME expert and reuse its weight pages. Chunking keeps the access pattern and leaves the kernel body untouched. Chunks are whole tokens so the kernel's own token = blockIdx.z / topk arithmetic stays valid against the offset vy / dst / topk_ids pointers. The 19 near-identical launcher bodies now share one moe_vec_launch<> helper. Tests assert the launch succeeds past the ceiling AND that rows below it are bit-identical however many chunks were used. The second is the real check: a wrong per-chunk offset would still launch, look healthy, and emit fluent nonsense. Verified to z = 160000 (20k tokens x top_k 8), and the kernel compiles for sm_75 and sm_89.
Ported from qwen3_5_moe/gguf_experts.py; llama.cpp emits the same three stacked tensors for both architectures, so only the arithmetic differs. DeepSeek-V4-Flash is 43 served layers of 256 experts at moe_inter_dim 2048 over dim 4096. Not yet reachable: the deepseek4 architecture is not registered and gguf.py does not exist, so nothing calls this. Landing it separately keeps the verified piece reviewable on its own. Verified against the real 164GB antirez/deepseek-v4-gguf Q4KExperts checkpoint: 43 layers uniform Q4_K on both banks, gate_up (256, 4096, 2304) and down (256, 4096, 1152), each cross-checked against the file's own tensor dims (Q4_K is 144 bytes per 256 elements, so H=4096 -> 16*144 and I=2048 -> 8*144). Two things carried over deliberately. The gate_up fusion reshapes to [E, I, rb] and concatenates on dim=1 so each expert keeps its own gate rows followed by its own up rows; cat(dim=0) then reshape is the version that looks right and lays every expert's gate down before any up, which loads and runs at full speed and emits fluent nonsense. And layers >= num_layers are skipped: the file carries a trailing NextN/MTP block, and counting it makes a uniform checkpoint look mixed.
parse_config recovers DeepseekV4Args from the checkpoint's inference/config.json, which ships beside safetensors. A standalone .gguf has no such file, so this rebuilds the same dataclass from KV keys. Nothing falls back to a dataclass default: a silently-defaulted hyperparameter here yields a model that loads and generates confidently wrong text. block_count does NOT relate to the MTP block consistently across architectures, so the served layer count is derived rather than assumed. qwen35moe counts its NextN block inside block_count (Ornith: 41 blocks, blk.40 IS the MTP block, 40 served); deepseek4 does not (43 blocks, blk.0..blk.42 all served, MTP carries no blk tensors). Subtracting nextn_predict_layers unconditionally silently drops the last real layer, which is exactly what the first version did. compress_ratios is the authority instead: one entry per served layer plus MTP. Entries != 0 mark a layer with an attention compressor and == 4 one with the lightning indexer, so the schedule is checked against the tensor table on every load. That check caught the off-by-one immediately (40/20 instead of 41/21) rather than leaving it to surface as degraded output after a 145 GiB load. Verified against antirez/deepseek-v4-gguf Q4KExperts: 43 served layers, 41 compressors, 21 indexers, expert types (Q4_K, Q4_K), score_func sqrtsoftplus from expert_gating_func 4, route_scale 1.5. iter_gguf_weights and convert_deepseek4_to_gguf still to come; the arch is not registered yet, so none of this is reachable.
Completes the adapter and registers the architecture. Reconciles against the real 164GB antirez/deepseek-v4-gguf Q4KExperts file: 1199 model parameters, 1199 matched, 0 unknown destinations, 0 shape mismatches, 0 unfilled. Four things did not carry over from the qwen adapters: block_count does not relate to the MTP block consistently. qwen35moe counts its NextN block inside it, deepseek4 does not. Subtracting nextn_predict_layers unconditionally dropped the last real layer; the compressor/indexer schedule cross-check caught it as 40/20 instead of 41/21. deepseek_v4 is raw nn.Module and loads via load_state_dict over named_parameters(), unlike the BaseOP-based qwen trees. FreeToken's GGUFLinear holds qweight as a plain attribute, which that loader cannot see, and assigning a non-Module over a Module child raises. Hence GGUFLinearNN / GGUFEmbeddingNN here, carrying the packed bytes in a non-grad uint8 Parameter named .weight so the loader's .to(p.dtype) cast is a no-op. dequant.dequantize's pure-torch path implements only Q4_0 and Q6_K; this checkpoint's attention projections and lm_head are Q8_0, which needs the CUDA kernel. Unquantized F16/F32 are reinterpreted rather than dequantized. Three tensors do not behave as their names suggest, and each would have been a silent mis-wire: attn_output_a is Q8_0 in the file but attn.wo_a is a bare bf16 nn.Parameter (no .weight, never packed); the compressor and indexer projections are F16, i.e. in GGML_UNQUANTIZED, so GGUFLinear's uint8 buffer cannot hold them; and Indexer.wq_b is declared Linear(kind="fp8") with a .scale no GGUF tensor can fill, so it is rebuilt as bf16 rather than populated. tid2eid is a routing index, not a weight: read as int32 and widened, never sent through a float dequant that would round token ids. Registration covers all five points: GGUF_ARCH_TO_REGISTRY, the register.py ModelSpec, _TOKENIZER_ARCH and _STOP_TOKENS, and the package exports. Stop tokens were read from the checkpoint's own vocab (eos id 1 <|end_of_sentence|>, plus <|EOT|> 128805) rather than assumed. Not yet loaded end to end: that needs ~145 GiB of pinned host RAM for the expert banks.
…lues GGUFLinear and GGUFEmbedding store every ggml type in a uint8 buffer of row_bytes width, including F32/F16/BF16, where "packed" just means the raw value bytes. fused_mul_mat_gguf's unquantized branch multiplied that byte view directly, so in_features came out as row_bytes -- twice too wide for F16 -- and the call failed with "mat1 and mat2 shapes cannot be multiplied (4x2048 and 4096x248320)". Only reachable when a checkpoint stores a projection unquantized, which is why it had not surfaced: Ornith's output.weight is Q6_K and takes a dequant path. Apodex-1.1-mini ships output.weight as F16 and fails at the first forward. The bytes are now reinterpreted through the type's real dtype before the matmul. The ACTIVATION is cast rather than the weight: converting the weight would copy the whole matrix per call (about 1 GB per forward for a 248k-vocab lm_head) and allocating that during CUDA graph capture fails outright, which is what the first version of this fix did. x is [tokens, hidden], so casting it is negligible, and computing in the stored precision matches what llama.cpp does for these tensors. Verified on Apodex-1.1-mini Q2_K (qwen35moe, 40 layers, 256 experts, F16 output.weight) on an 8GB RTX 4060: serves, and correct at temperature 0 -- "The capital city of France is" -> " Paris", "The largest planet in our solar system is" -> " Jupiter. Its diameter is approximately 142,98[4 km]".
nekomario28
marked this pull request as ready for review
August 26, 2026 05:26
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Internal integration carrier only. Base is current ROCm shadow
cce26fac70661f1e8e0e9f8d3a16d7e511654c9b; head is exact upstream FlashML-org#131 donor7da6b77e3a20eaa310d246650745b05538661b5a. Purpose: preserve donor history, expose merge conflicts, and build the smallest ROCm × qwen35moe/GGUF validation branch. Not an upstream proposal and not intended for fork main.