GGUF: support all quant types, add qwen35moe

[vcruz305]

Two related things here: GGUF quant types beyond the three the Python side knew about, and qwen35moe support so llama.cpp MoE checkpoints load. I did the second one first and the first one fell out of it, so they share a branch. Happy to split if you'd rather review them separately.

Tested on an RTX 4060 Laptop (8GB) with Ornith-1.5-35B-A3B IQ3_S, a 16GB qwen35moe quant, experts offloaded to host RAM.

Generic quant types

csrc/gguf/ already dispatches 19 ggml types but the Python tables only described 6, so K-quants and I-quants were unreachable. BLOCK_SHAPE now covers all 21 types and there are five frozensets mirroring the C dispatch switches, each with the file:line it mirrors so they can be re-checked. tests/models/test_gguf_type_tables.py derives the block sizes from ggml-common.h at runtime and extracts the case labels from gguf_kernel.cu, so the tables can't drift from the kernels without a test failing.

One thing worth flagging on its own: none of the five switch (type) blocks in gguf_kernel.cu had a default:. Y is allocated with torch::empty, so an unsupported type returned uninitialized memory instead of raising, and ggml_moe_get_block_size returned 0. They all raise now. That's a latent bug independent of everything else in this PR.

I-quants have MMVQ and dequant kernels but no MMQ kernel, so fused_mul_mat_gguf routes them through ggml_dequantize plus a torch matmul at prefill. That branch already existed but was unreachable, since _MMQ was tested before _DEQUANT and the two sets were identical.

qwen35moe

New models/qwen3_5_moe/gguf.py and gguf_experts.py, following the gemma4 pattern. The config comes from GGUF metadata and the head layout is derived rather than assumed: attn_qkv's packed width of 8192 is exactly q(group_count * state_size) + k(group_count * state_size) + v(time_step_rank * state_size), and parse_gguf_config asserts the ssm.inner_size relation instead of trusting it.

Four things llama.cpp's converter does that the loader has to undo. I found these by reading conversion/qwen.py after a lot of wasted time, so listing them in case they save someone else the trouble:

• ssm_a holds A = -exp(A_log), not A_log. Every value in the file is negative, spanning roughly -70 to -0.02. gdn.py computes -A_log.exp(), so passing A straight through makes the recurrent decay gate about -2.5e-5 and the recurrence dies. The loader yields log(-A).
• The (1 + w) RMSNorm shift is already folded in by the converter. weight.py adds it for HF checkpoints because HF stores raw w, but a GGUF arrives pre-shifted and adding it again double counts. The measured means gave this away, around 1.0 rather than the 0.0X raw w would give.
• V heads are stored tiled rather than grouped when num_k_heads != num_v_heads, so ggml_repeat can replace an interleaved repeat. FreeToken pairs K head k with V heads [k*R, (k+1)*R) like HF does, so all seven affected tensors get un-tiled on load. Row permutations are done on the packed bytes directly, which is safe because each output row is its own run of blocks. ssm_out is the exception since its columns carry the V dimension and a 128 wide head straddles the 256 element blocks, so that one tensor is dequantized to dense bf16, about 500MB across 30 layers.
• Merged projections are not uniformly typed. In IQ3_S, qkv_proj is q and k IQ3_S with v Q4_K, and the GDN in_proj is attn_qkv Q4_K with the rest IQ3_S. Concatenating packed rows needs matching row_bytes (880 vs 1152 here), so there's a GGUFMergedLinear that keeps one buffer per part and concatenates the outputs instead. Same result as the merged GEMM since every part reads the same input. gguf_merged_or_plain() keeps the uniform case on the cheaper single kernel path.

None of these are visible to shape, dtype or byte-identity checks. The model loads, runs at full speed and produces fluent nonsense, which is a slow way to find out.

Expert banks

_BANK_SCHEMAS gets a "gguf" entry, and the (gate_up, down) ggml types are threaded from ModelConfig through ExpertBanks and OffloadMoeCache to the kernels, because a GGUF's row stride is a property of the file rather than the format. fused_experts_gguf takes a separate down_quant_type since the two banks have separate slot pools and can legitimately differ.

What is not supported is a bank whose 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 every layer, so two strides in one pool would read every block at the wrong offset. _gguf_banks raises and names the offending layers. In practice this means llama.cpp's mixed levels don't load: for Ornith-1.5-35B, IQ3_S and IQ3_XXS have uniform banks while IQ3_M, IQ2_M, IQ2_XXS and IQ1_S split ffn_down_exps across two types. Quantizing with --pure avoids it. Per type slot pools would fix it properly but that's surgery on the hot path and I didn't want to bundle it here.

_model_setup_override now declines for expert_quant == "gguf", otherwise qwen3_5_moe's NVFP4 setup_offload_expert_banks hijacks the GGUF load and fails deep inside an NVFP4 reader.

Tokenizer

convert_gguf_tokenizer is keyed on transformers' own model_type, so passing the GGUF arch string through raised KeyError('qwen35moe') and killed the detokenizer worker during load. qwen35moe is a GPT2 style BPE so the existing qwen2 converter handles it. Two neighbouring gemma4 assumptions are now per-arch: eos preferred the literal <turn|> and the stop ids were hardcoded to <eos>/<turn|>, neither of which exists in a Qwen vocab, and unk defaulted to a <unk> token that also isn't there. gemma4 behaviour is unchanged.

Numbers

load                65s (serial GGUF read)
VRAM                6879 / 8188 MiB
host RAM            20GB pinned expert banks
decode              46.7 to 50.1 tok/s
end to end          29.4 tok/s including prefill
llama.cpp CPU ref   11.1 tok/s

Correctness checks I ran, since "it produces text" isn't much of a claim:

• 573 weights reconciled against the constructed module: every name exists, every shape and dtype matches, no non-expert buffer left unfilled.
• Expert banks checked byte for byte against their source rows, per expert, at both ends and the middle. This caught the gate_up fusion laying down every expert's gate before any up.
• MMVQ, MMQ and the grouped MoE GEMV checked against dequant plus matmul on real tensors from the file, cos 0.9999 or better for IQ3_S, Q4_K and Q6_K.
• 8/8 on short factual questions with greedy decoding, including arithmetic.

First token argmax matches llama.cpp CPU on 2 of 6 single token prompts. Every disagreement is a plausible near tie (Paris continuing " de" against " is", import continuing " java" against " sys"). FreeToken runs CUDA W4A8 MMVQ and llama.cpp runs CPU AVX, so I don't think exact agreement is reachable, but I'd rather state the number than imply it's bit exact.

Known gaps

• TP=1 only, same as the gemma4 GGUF path.
• The NextN/MTP block is dropped, so no speculative decoding.
• Bank loading is serial. _gguf_banks raises NotImplementedError for the parallel reader since that path assumes safetensors and a GGUF is one packed file.
• max_model_len still reports the metadata context length even when KV is capped lower.
• I only tested the flashinfer attention backend. The triton fallback is untouched but unexercised here.

[vcruz305]

Thank you for this. Reproduced, and it is fixed in b2f8475.

Your diagnosis is right and your read of the consequence is better than mine was. I had found the same missing slice about an hour before you posted, but I came at it from the memory side: on an 8GB card a ~1800-token prompt was killing the worker, because logits are [tokens, vocab] and Ornith's vocab is 248,320, so that is 486 KiB of logits per token and a single 894 MB allocation. I fixed it as a memory bug and moved on.

I missed that forward_batch does logits[: batch.size], so on smaller prompts that never OOM you do not crash, you just sample the first token off position 0 and get a quietly wrong answer. That is the worse half of the bug and you caught it. The Hello! The capital city of France is case is a good repro precisely because it fits comfortably in VRAM, so nothing looks wrong.

I went with the same slice you proposed, in a slightly different place. Rather than a nested class in convert_qwen35_to_gguf, it is GGUFLMHead in layers/gguf.py:

class GGUFLMHead(GGUFLinear):
    def forward(self, x: torch.Tensor) -> torch.Tensor:
        from freetoken.core import get_global_ctx

        batch = get_global_ctx().batch
        if batch.is_prefill:
            indices = batch.attn_metadata.get_last_indices(batch.size)
            x = x[indices].contiguous()
        return super().forward(x)

Two reasons for the move. qwen3_moe/gguf.py has the identical untied branch and the identical bug, so it needed fixing in both places. And the slice deliberately does not go into GGUFLinear itself, since that class is also every q/k/v/o and MLP projection in the model and must not slice.

I checked the other adapters: gemma4 always takes the tied path, so GGUFTiedLMHead covers it, and those two were the only untied sites.

Verified on Ornith-1.5-35B-A3B IQ3_S, untied lm_head, temperature 0:

Hello! The capital city of France is  ->  " Paris. Paris is located in the northern part of France"
The capital of Japan is               ->  " Tokyo, a city that has been the seat of the"

and on the memory side, same box, 8GB, offload backend: ~1,800 tokens went from killing the worker to 6.1s, and 4,598 tokens now prefills in 16.9s.

One thing you may hit next on an 8GB card. Past roughly 5-6k prompt tokens it still runs out of room, but somewhere else: the GDN chunked kernel's activations, chunk_gated_delta_rule_fwd -> chunk_fwd_o. That one is headroom rather than a defect, and --max-prefill-length 2048 handles it (6,160 tokens in 23.2s, 7,784 in 27.5s on the same box). Worth knowing that the 8192 default does not chunk a typical long prompt at all, so you get one enormous forward pass.

No need to send a PR for this one, but the report was genuinely useful and I would take more of them.

[max-nuv]

I really hope you can get this fix working on Windows, because after pulling your branch and fixing the build for MSVS, I got a CUDA error: invalid argument in ggml_moe_a8_vec with the first 20k tokens prompt. After running short tests on Qwen 3.6 35B A3B (NVFP4) vs. Ornith 1.5 35B A3B (IQ3_XXS), I got the following results:

Test	Qwen	Ornith
Logic TTFT	1.09 sec	0.78 sec
Logic speed	26.7 tokens/sec	40.9 tokens/sec
Logic quality	Hit the limit.	Correct: $100, finish=stop
Code TTFT	0.89 sec	0.69 sec
Code speed	27.4 tokens/sec	40.0 tokens/sec
Code quality	Hit the limit.	Correct code, but finish=length

So, Ornith, with a smaller quantum, produced better and faster results, and I'm really looking forward to trying the model in real-world tasks.

P.S. 3070 8GB VRAM, context 96k tokens

[vcruz305]

Found it and fixed it in 9952a39. Thank you for the report, and for building it on MSVC to get there.

It is not a Windows problem and not a 3070 problem, which is worth saying because it would have sent me looking in the wrong place. moe_vec.cuh carries the routed-pair index (token * top_k + slot) in gridDim.z, and CUDA caps gridDim.z at 65535 on every architecture. All 19 launchers computed:

const dim3 block_nums(block_num_y, 1, tokens * top_k);

At top_k 8 that ceiling is 8191 tokens. Your 20k prompt asks for z = 160000, the launch is rejected with cudaErrorInvalidConfiguration, and torch surfaces it as the fairly unhelpful "invalid argument". Anyone with a long enough prompt hits it on any GPU and any OS.

fused_experts_gguf makes two launches, gate_up as (top_k, N) and down as (1, N * top_k), and both land on z = N * top_k, so they overflow together.

The launcher now issues token-aligned chunks. Moving the pair index into x would also have fixed it, 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 preserves that 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 duplicated launcher bodies now share one moe_vec_launch<> helper.

Tests in tests/kernels/test_moe_vec_grid_z.py check two things, and the second is the one that matters:

• a launch past the ceiling succeeds (verified to z = 160000, i.e. 20k tokens at top_k 8)
• rows below the ceiling are bit-identical however many chunks were used

A wrong per-chunk pointer offset would still launch and still look healthy, and would quietly corrupt everything above 8191 tokens. The equality check is what catches that; a crash test would not.

Also verified the kernel still compiles for sm_75 and sm_89.

On your numbers: Ornith IQ3_XXS beating Qwen NVFP4 on both speed and quality is a real effect rather than a fluke of the quant. Ornith's expert banks come out uniform at IQ3_XXS, so the whole MoE stays on the packed moe_vec path with no dequantization, whereas the NVFP4 route materializes more per token. The finish=length on both Qwen tests also suggests it was running into the output cap rather than answering worse, so the quality column may be flattering Ornith slightly.

One thing you should hit next, at 96k context. There are two other prefill limits I have measured, neither of which is this bug:

• An untied lm_head was projecting every prefill position instead of the last, so logits were [tokens, vocab]. Fixed in b2f8475. On a large-vocab model that is hundreds of MB per prompt and it also made the first sampled token come from position 0.
• Past roughly 5-6k prompt tokens on an 8GB card the GDN chunked kernel's activations run out of room (chunk_gated_delta_rule_fwd -> chunk_fwd_o). That one is genuine headroom rather than a defect, and --max-prefill-length 2048 handles it. Worth knowing the 8192 default does not chunk a typical long prompt at all, so you get one enormous forward pass.

If you pull again you will get all three. I would be glad of more reports like this one.

[spacewave-bot]

Verified this end-to-end on an AMD Radeon RX 7700 XT (gfx1101, ROCm 7.2/HIP), using the HIP engine port from #137 plus this patch set: Qwen3-30B-A3B IQ4_XS serves coherently with --moe-backend offload at approximately 52 tokens/s decode (eager mode; CUDA graph capture currently requires --cuda-graph-max-bs 0, see below).

Three small HIP fixes were required on top, submitted as vcruz305#1:

1. gguf.py passes nvcc-only flags (--expt-relaxed-constexpr, -ccbin) to the HIP build; hipcc rejects both, and relaxed constexpr is already its default.
2. The warp-sync calls in mmvq.cuh, moe_vec.cuh, and gguf_kernel.cu pass 32-bit masks; HIP requires 64-bit. CUDA accepts the 64-bit form as well, so the change is unconditional.
3. layers/activation.py routes GGUF-MoE activations through a triton kernel that fails LLVM register allocation (couldn't allocate output register for constraint 'f') on some RDNA3 shapes; under HIP this falls back to plain torch ops.

Also worth noting for other testers: setting PYTORCH_ROCM_ARCH=<your gfx> reduces the initial GGUF kernel build time substantially.

[zyy0212time-del]

Windows/CUDA follow-up: end-to-end verified + a request on merge timing / Windows wheels

Thanks again for b2f8475 and 9952a39. Reporting back from the Windows side for completeness.

Our engine is this branch @ g30aa89115 (it predates both fix commits) with equivalent local patches for the two bugs, so this validates the diagnosis and approach rather than the exact commits:

• Qwen3.6-35B-A3B Q4_K (pure-quant, abliterated) on an 8GB RTX 5060 Laptop, --moe-backend cpu: ~70s cold start, ~13.4 t/s decode, chat output correct.
• Prompts around 8K tokens (agent frameworks with large system prompts) no longer crash the worker - before the chunking fix they reliably took down ggml_moe_a8_vec.

One practical concern: FreeToken Desktop on Windows keeps showing the "engine update available" prompt. The latest release (v0.1.2, published 2026-08-19, before this PR was opened) ships Linux-only wheels and has no qwen35moe GGUF support, so Windows users depending on this branch currently have no safe update path.

Is there a timeline for merging this? And could the next release include win_amd64 wheels? Happy to help test Windows builds (CUDA 13.0 / MSVC 17.14) - this machine is set up and idle.