diff --git a/.agents/claims/CLAIM-QUANT-EXL3-DOCS-IMPLEMENT.md b/.agents/claims/CLAIM-QUANT-EXL3-DOCS-IMPLEMENT.md new file mode 100644 index 000000000..11e11e5b0 --- /dev/null +++ b/.agents/claims/CLAIM-QUANT-EXL3-DOCS-IMPLEMENT.md @@ -0,0 +1,5 @@ +# CLAIM-QUANT-EXL3-DOCS-IMPLEMENT + +| Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update | +|---|---|---|---|---|---|---|---| +| `CLAIM-QUANT-EXL3-DOCS-IMPLEMENT` | `QUANT-EXL3` (`ACTIVE`) | Codex (gpt-6), helper for the documentation spec | `/home/mudler/_git/vllm.cpp-wt-exl3-docs-implement`; CPU-only | `row/QUANT-EXL3-DOCS-IMPLEMENT` | README old EXL3 news; FEATURES EXL3 quantization row; USAGE EXL3 artifact rows; dated prose archive; `specs/exl3-public-docs.md` outcome. No implementation, GPU, or other claim's files. | `GATING` | 2026-09-13: implementation ready for independent review in PR #3176; scoped editorial work for `ISSUE-LOCAL-01M2CBGG169HDX0DWKHQWCBBGN`; parent owns integration and broad preflight | diff --git a/.agents/completed/exl3-public-docs-20260913.md b/.agents/completed/exl3-public-docs-20260913.md new file mode 100644 index 000000000..c13c3e2fd --- /dev/null +++ b/.agents/completed/exl3-public-docs-20260913.md @@ -0,0 +1,48 @@ +# EXL3 public prose before the 13 September 2026 correction + +Frozen excerpts from `0cf30baae`, preserved verbatim for historical evidence. +These excerpts contain superseded claims; use the live EXL3 specs for current decisions. + +## README.md + +```text +- **2026-08** **EXL3 checkpoints now generate on CPU and CUDA.** A stock + Llama-3.2-1B-Instruct EXL3 checkpoint loads through the shared dense model path and emits text. + The current CUDA path supports its 3-bit body and 6-bit output head. No speed claim is available. +``` + +## docs/FEATURES.md + +```text +| EXL3 trellis (exllamav3; codebooks 0 (3INST), 1 (MCG) and 2 (`mul1`), Hadamard-128 + sign vectors, NO scales) | ◐ **A stock EXL3 checkpoint GENERATES**: `turboderp/Llama-3.2-1B-Instruct-exl3` @ 3.0bpw loads through the shared dense container and emits coherent text from `vllm-cli` on a CUDA queue and on a CPU queue, which reaches `LlamaForCausalLM` and Qwen3-dense together. The scheme sits on vLLM's own `LinearMethodBase` seam, so it is no longer a DeepSeek-V4-private arm. `bits` and the codebook are both read PER TENSOR — the published 3.0bpw Llama has a 3-bit body and a 6-bit head, and ships no `mcg` marker, which means codebook 0 and not MCG. **The DeepSeek-V4 arm is separate and unchanged**: its rank-sliced SparkInfer artifact loads and executes end to end on a synthetic checkpoint. The rank-sliced `0xSero/deepseek-v4-flash-0731-spark` routed-expert tower coalesces TP4->TP1 at load, the `carried-*` half (block-wise FP8 + BF16 + F32) is dequantized into the host-float tower the forward composes with, and `DeepseekV4Model::Forward` runs the whole model, dispatching one `vt::Exl3MoeMlp` per MoE layer over the routed experts (`bits == 3`, `mcg` codebook). **The REAL artifact still does not run**: its DSA compressor and indexer tensors are twice the width the host forward indexes and the loader refuses them by name, and its tokenizer is not read (#1924). **The device half is now PARTLY VERIFIED** (GB10 `sm_121a`, 2026-08-28): `had_r_128` is BYTE-IDENTICAL CUDA-vs-CPU, `exl3_gemm` matches the f64 reference at `rel_rms 5.538e-4` against a `1.0e-3` bound, and the `m<=8` GEMV meets tier 3c at `5.160e-4` against `6.0e-3`. What is STILL unverified on a device is the FUSED MoE arm, which cannot run on this code at all: it needs a device-resident tower and `CudaBackend::DeviceMemoryIsHostAddressable()` is false by design, so the routed-expert path executes on a CPU queue today. No speed number is claimed on any axis. **Codebook 2 (`mul1`) and the 4- and 5-bit widths are now implemented** so that `Mia-AiLab/Qwen3.8-27B-EXL3-3.5bpw` -- **409** trellis modules, 137 at 3 bpw, 270 at 4, one at 5 and one at 6, every one of them `mul1`-marked -- is no longer refused by name ([#2495](https://github.com/mudler/vllm.cpp/issues/2495)). The count read 272 until 2026-09-02 and was an UNDERCOUNT that omitted the 137 bits-3 modules, which is why `(3, 2)` was missing from the CUDA instantiation list and every MLP projection of that target refused on a CUDA queue. **`(3, 2)` is now instantiated and gated on `dgx:gpu0`** (GB10, driver `580.173.02`), agreeing with the CPU decoder at `rel_rms 3.02544e-07` against a `1.0e-3` bound ([#2574](https://github.com/mudler/vllm.cpp/issues/2574)). cb 2 is a DIFFERENT decode and not a third multiplier: it sums the four bytes of the product into an fp16 bit pattern and maps it with a fused fp16 affine, and the host arm is gated against hand-computed upstream values rather than against itself. **A ROCm arm now exists too** ([#2433](https://github.com/mudler/vllm.cpp/issues/2433), [spec](../.agents/specs/backend-rocm-exl3.md)): `Exl3Gemm` on `gfx1151` is a transcription of the PORTABLE CPU reference rather than a port of the CUDA kernel -- 90 KiB of requested shared memory does not fit a 64 KiB AMD LDS, and `m16n8k16`, `ldmatrix` and `cp.async` have no AMD form -- and because it uses no matrix cores it is gated at BYTE equality with the CPU arm on all three codebooks, not at a tolerance. It covers every width the host arm decodes, so the CUDA instantiation list below does not bound it. On the CUDA device arm `(3,0)`, `(3,1)`, `(6,0)`, `(3,2)`, `(4,2)`, `(5,2)` and `(6,2)` are instantiated -- `(3,2)` is the width the undercount cost, added and device-gated by [#2574](https://github.com/mudler/vllm.cpp/issues/2574); every other width and codebook and the device-resident tower are owed ([spec](../.agents/specs/quant-exl3-mul1.md)). **The CUDA GEMM now has a dual-path dispatch mirroring exllamav3** ([#3124](https://github.com/mudler/vllm.cpp/issues/3124), [spec](../.agents/specs/quant-exl3-shared.md)): M <= 144 routes to the cooperative `exl3_gemm` kernel; M > 144 routes to `Exl3ReconstructGemm`, which dequantizes trellis weights to fp16 on-device and runs cuBLASLt fp16 GEMM. Only CUDA registers the reconstruct kernel, so CPU, ROCm and Vulkan keep `exl3_gemm` at every M, as they did before the dispatch existed. The fused path (M >= 1024, both dims 128-divisible) folds the input/output Hadamards into the reconstruct kernel; the unfused path applies them separately. N > 32768 slices the output dimension. All 7 (bits, codebook) arms are gated against an f64 reference at `rel_rms <= 1e-3` on `dgx:gpu0`. **The `m<=8` GEMV is a SEPARATE arm set from that list and a separate row** ([`QUANT-EXL3-PERF`](../.agents/specs/quant-exl3-perf.md), [#2570](https://github.com/mudler/vllm.cpp/issues/2570)): upstream instantiates `(4,0) (4,1) (4,2) (2,1) (2,2) (3,1) (3,2)` and this tree now carries `(3,1)`, `(3,2)` and `(4,2)`, so **407 of that checkpoint's 409 trellis modules** reach the small-m fast path where before NO module of it could -- the only arm was `(3,1)` and the artifact contains zero `(3,1)` tensors. `(3,2)` was one template argument; `(4,2)`, the 270-module arm, was a KERNEL port: a per-width `LSTRIDE`, two 24-lane guards that exist only because a 3-bit tile is 24 uint32 rather than 32, and a separate `dq8_regs_4bits` window read. `(4,0)`, `(4,1)` and the whole 2 bpw arm stay OWED -- no artifact in this tree carries a 4-bit tensor at either of the first two codebooks, and nothing builds them, so that the kernel would compile for them is an inference and not a measurement. **Instantiating an arm is necessary and not sufficient**: `Exl3GemvSelectConfig` is upstream's envelope verbatim and DECLINES, and on Blackwell EVERY shape of this checkpoint, at both widths, is admitted only where `size_n / 32 <= narrow_coresident`, an occupancy query. The wide config is not an escape from it: its band needs `size_k <= 4096` and the smallest 4-bit `k` here is 5120. The thresholds are `>= 544` and `>= 160` at the two bits-3 shapes and 32, 160, 192, 320, 384 and 544 across the six bits-4 ones, all pinned in `tests/vt/test_exl3_gemv.cpp` from both sides. **The throughput effect is now MEASURED on GB10, and it is a NULL** ([#2570](https://github.com/mudler/vllm.cpp/issues/2570)): three interleaved rounds on one binary read 17.20/17.18/17.09 tok/s with the arm off against 17.22/17.12/17.09 with it on, a 1.3% spread and no separation. `nsys --cuda-graph-trace=node` shows why, and it is a DECLINE and not an ineffective kernel: `exl3_gemv_kernel` appears 0 times at the default and 4 times only when forced. GB10 reports `SM_COUNT=48 MAX_THREADS_PER_SM=1536`, so the narrow config's 512-thread blocks ceiling `narrow_coresident` at 144, and only shapes needing `<= 144` are admitted -- 34 of the checkpoint's 409 modules, 0.75% of its trellis bytes. The arms are correct, upstream-faithful and device-gated; on THIS part the envelope declines them. `m == 8` is unmeasured, because the draft legs did not run. **The FUSED MoE mgemm is a SEPARATE arm set from that list**, and upstream's own bound on it is narrower: `exl3_moe.cu:184` admits `mcg` and `mul1` and refuses 3INST, so codebook 0 is not a MoE arm upstream either. It now instantiates bits 3, 4, 5 and 6 at codebook 1 where it carried bits 3 alone ([#2756](https://github.com/mudler/vllm.cpp/issues/2756)); before that a mcg expert tower at any other width could not run on a CUDA queue by ANY path, because the fused arm is default-ON and its refusal is not caught, and the `VT_DSV4_EXL3_FUSED_MOE=0` rollback lands on a GEMM arm set with no `(4,1)`, `(5,1)` or `(6,1)` either. Codebook 2 on the fused arm is owed as a LOADER slice and not a kernel one: five sites pin the codebook to 1 before the kernel is chosen, so a cb-2 instantiation added alone would be dead code. **Its GDN linear-attention tower now loads and runs too** ([#2495](https://github.com/mudler/vllm.cpp/issues/2495) item 4, [spec](../.agents/specs/model-qwen35-gdn-exl3.md)): 48 of that model's 64 layers are `linear_attention`, and the three projections of each were refused by name because nothing in the GDN forward consumed an `Exl3Weight`. The arm holds THREE trellises rather than one merged `in_proj_qkvz` owner, because the artifact ships `in_proj_qkv` and `in_proj_z` as two independently quantized tensors whose sign vectors are fitted per projection, so the bf16 merge has no trellis analogue. It also accepts the F16 `in_proj_a`/`in_proj_b` the artifact stores beside its BF16 `conv1d`, `norm`, `A_log` and `dt_bias`. **The DFlash2 DRAFT loader now has an EXL3 arm too** ([#2495](https://github.com/mudler/vllm.cpp/issues/2495) item 7, [spec](../.agents/specs/model-dflash2-exl3.md)): `Mia-AiLab/Qwen3.8-27B-DFlash2-EXL3-5.0bpw` ships all 36 of its modules at bits 5 with a `mul1` marker, and `LoadQwen3DFlash` died on `fc.weight`, which an EXL3 module does not ship. The arm holds SEVEN trellises per layer rather than the merged `qkv_proj` and `gate_up_proj` owners the bf16 reader builds, for the same reason the GDN tower holds three. It also accepts the F16 `candidate_selector.hidden_projection` and `*_conv.kernel_projection` the repack left unquantized beside their BF16 `base_kernel` and codebooks — a second refusal that was invisible until the first was gone. **Its `mtp.*` draft head loads and runs too, and all THREE readers of the target's trellis head now compute with it packed** ([#2495](https://github.com/mudler/vllm.cpp/issues/2495) items 5 and 6, [spec](../.agents/specs/model-qwen35-exl3-head.md)): the MTP loader refused nine quantized `mtp.*` tensors as "expected BF16", and the MTP draft's `ComputeLogits` plus the DFlash/DSpark draft's SHARED head read knew only bf16 and NVFP4. The head is COMPUTED WITH rather than widened, which is what upstream does without a branch because its head is an `nn.Module` and `_apply_head` calls `lm_head.quant_method.apply`; a dequantized copy of the real 248320x5120 head would be 2.543 GB against 0.953 GB packed. The shared read also used to fall off the end of its bf16 loop and return SILENTLY with every owner empty when no arm matched, which is exactly what this target does -- it ships no `lm_head.weight` at all -- and it now refuses by name ([#2569](https://github.com/mudler/vllm.cpp/issues/2569)). **The checkpoint's real inventory is 409 quantized modules, not the 272 this row previously recorded**, measured from both downloaded shards' headers: bits 3 x137, 4 x270, 5 x1, 6 x1, every one `mul1`, and no tensor in the file has `scale` in its name. `(bits 3, codebook 2)` is now instantiated on the CUDA GEMM and gated on a device, so those 137 modules no longer refuse by name ([#2574](https://github.com/mudler/vllm.cpp/issues/2574), owned by `QUANT-EXL3-MUL1`, not by this row). **THE 27B CHECKPOINT NOW RUNS ON CUDA AND GENERATES.** On `dgx:gpu0` (GB10, driver 580.173.02, `sm_121a`) the published `Mia-AiLab/Qwen3.8-27B-EXL3-3.5bpw` loads and emits ` Paris. The capital of Germany is Berlin. The capital of Italy is` at `rc=0`, with both shard sha256 values recomputed on the device. Warm decode is **16.7 tok/s** target-only and **48.7 tok/s** with its DFlash2 EXL3 draft at k=7, interleaved on one binary, and the two arms are **token-identical** ([`qwen38-27b-exl3-gb10`](benchmarks/qwen38-27b-exl3-gb10.md)). **That is NOT a reproduction of the upstream README's 47.5 tok/s**, and must not be quoted beside it: theirs is HumanEval-style at T=0.6 with acceptance 4.43, this is greedy T=0 on a prompt that continues into a list of capitals, the context is 8192 rather than 262144 with no NVFP4 KV cache, and the shipped paged draft route had to be disabled (#2274). A matched HumanEval-style run at T=0.6 with acceptance reported is owed | ☐ no EXL3 at the parity pin | ☐ | ☐ | +``` + +## docs/USAGE.md + +```text +| Qwen3.8-27B EXL3 3.5bpw (the #2495 benchmark target; RUNS and GENERATES on CUDA) | `model-0000{1,2}-of-00002.safetensors` | 15,338,408,461 bytes total (index `total_size` 15,338,106,948 plus the two safetensors headers); 2426 tensors | `Mia-AiLab/Qwen3.8-27B-EXL3-3.5bpw` @ `19441ac874c4018295da848e250f23511361cda4` | `7b77214fe58ff15fed0b4af55e3cd92f38842b8711886d68954e8071ff8270c6` (shard 1), `411c83bb1070b27f3d670fc93e38dca0f17eb66429f64b5706901b12613188b2` (shard 2) — both **recomputed on `dgx:gpu0`** and matching the download host | `Qwen3_5ForConditionalGeneration`, 64 layers as 48 `linear_attention` + 16 `full_attention`. **409 trellis modules, every one codebook 2 (`mul1`), zero `mcg`**: bits 3 x137 (`mlp.{gate,up,down}_proj` of 46 layers), bits 4 x270 (all 48 GDN `linear_attn` x3, 55 more `mlp.*`, the 16 `self_attn.*`), bits 5 x1, bits 6 x1 (`lm_head`). It also ships its own quantized MTP head, 39 tensors under `mtp.*`, uniformly bits 4. GENERATES on GB10 greedy at `rc=0`; 16.7 tok/s target-only and 48.7 tok/s with its DFlash2 draft ([`qwen38-27b-exl3-gb10`](benchmarks/qwen38-27b-exl3-gb10.md)) | Its vision tower (`language_model_only: false`, depth 27) is not exercised. `quantization_config` says `out_scales: "always"` but the artifact ships **no scale tensors** — every quantized module is exactly `{trellis, suh, svh, mul1}` and the setting is folded into `svh`. The `m <= 8` GEMV had an arm for **none** of its 409 modules until 2026-09-02 -- its only arm was `(3,1)` and this artifact has zero `(3,1)` tensors. `(3,2)` is now instantiated, which is the 137 bits-3 modules, and `(4,2)` too, which is the other 270 ([#2570](https://github.com/mudler/vllm.cpp/issues/2570), [`QUANT-EXL3-PERF`](../.agents/specs/quant-exl3-perf.md)); 407 of the 409 now have an arm, and the 5- and 6-bit pair have none upstream either. **Whether the arm is TAKEN is a separate question from whether it exists**: `Exl3GemvSelectConfig` is upstream's envelope verbatim and declines on Blackwell below an occupancy threshold -- `narrow_coresident >= 544` at k 5120 n 17408 and `>= 160` at k 17408 n 5120, and 32, 160, 192, 320, 384 and 544 across the six bits-4 shapes -- and that HAS now been measured on GB10: `SM_COUNT=48 MAX_THREADS_PER_SM=1536` ceilings `narrow_coresident` at 144, so 34 of the 409 modules are admitted and the decode A/B is a null (17.16 vs 17.14 tok/s mean over three interleaved rounds, 1.3% spread), with `nsys` showing zero `exl3_gemv_kernel` launches at the default. The wide config does not escape it here, because its band needs `size_k <= 4096` and the smallest 4-bit `k` in this artifact is 5120. No throughput change is claimed here | +``` + +## docs/USAGE.md + +```text +| Qwen3.8-27B DFlash2 draft, EXL3 5.0bpw (the drafter for the row above) | `model.safetensors` | 1,470,916,078 bytes; 189 tensors | `Mia-AiLab/Qwen3.8-27B-DFlash2-EXL3-5.0bpw` @ `4f0436269bca761b071f05319e8e04a87cc633f9` | `6b2e3afc694a343b7f3f0edfe5925e460762fc9ede4699165b577ca0733c8e56` | `DFlash2DraftModel`, 5 sliding-attention layers, `block_size 8`, taps at target layers `[5,19,33,47,61]`. **36 trellis modules, all bits 5 codebook 2** — uniform, NOT the "module-adaptive" the upstream README claims. It owns no `lm_head` and no `embed_tokens`: both are SHARED from the target, and the head is consumed **packed** (2,542,796,800 B saved on device). Drafts on GB10 at k=7 with output token-identical to the target-only arm | The candidate selector is **not** quantized, contrary to that README — `candidate_selector.hidden_projection` is dense F16 and its two `[248320, 256]` codebooks are raw BF16, together ~17% of the file. Its `kernel_projection` and selector weights are **F16** where the surrounding norms are BF16, admitted by name only. The shipped **paged** draft route faults on run 2 in one process ([#2274](https://github.com/mudler/vllm.cpp/issues/2274)); `VT_DFLASH_PAGED=0` is the measured arm | +``` + +## docs/USAGE.md + +```text +| Llama-3.2-1B-Instruct EXL3 3.0bpw (the first EXL3 checkpoint that GENERATES) | `model.safetensors` | 1,089,087,416 bytes | `turboderp/Llama-3.2-1B-Instruct-exl3` @ `f8f438c290680b15622270eff03bef23a458b1cf` (revision `3.0bpw` -- this repo publishes ONE BRANCH PER BIT WIDTH and `main` carries no weights at all, so a bare repo id resolves to nothing) | `3c0341e9c7c4c16a86a499de1dff4f6d7de9855541d669f3b0e214d72b54c2fc` | LOADS and GENERATES end to end through `vllm-cli` on `--device cpu`: `The capital of France is` -> ` Paris. Paris is known for its famous landmarks such as the Eiffel Tower` (greedy, 16 tokens, 2026-08-28). Native exllamav3 layout, no `.rank{r}` slicing; the body is 3-bit and `lm_head` is SIX-bit, resolved per tensor | Codebook **0** (the original QTIP 3INST), because the artifact ships no `mcg` marker and `LinearEXL3` derives the codebook from tensor PRESENCE. **It now RUNS ON CUDA**: the device arm instantiates `(3,0)`, `(3,1)` and `(6,0)`, so this checkpoint's 3-bit body and 6-bit head both reach the GPU. **It now RUNS ON ROCm too**, on `strix:gpu0` (`gfx1151`, RDNA3.5 APU, ROCm 7.2.4, 2026-09-02): the same file completes greedy generation with ZERO CPU reference-tier operations at 8.27 tok/s warm, against 0.83 tok/s for the same tree with the two ROCm registrations disabled, and the two arms emit the IDENTICAL continuation ` Paris. Paris is known for its famous` ([#2433](https://github.com/mudler/vllm.cpp/issues/2433)). That is an EXL3-vs-EXL3 A/B and NOT a ratio against a bf16 target: the BF16 control hung the GPU in the same lease ([#2511](https://github.com/mudler/vllm.cpp/issues/2511)), so this row records no BF16 denominator and no AMD clock attribution. gfx1151 is an APU with unified memory, so the figure generalizes to no discrete Radeon -- and on a discrete board the reference tier is off by design, which means this arm is what makes EXL3 RUN there at all rather than what makes it fast. It has NO GEMV fast path at `m == 1` and takes the regular shape table, which is upstream's behaviour too — its envelope refuses `bits != 4 && cb == 0` and its instantiation list omits `(3,0)`. **No speed is claimed.** The device figure ranged 2.1-5.0 tok/s across jobs on one binary and box, and an interleaved comparison against the BF16 twin of the same model read 2.33 vs 2.24 tok/s — indistinguishable, so the remaining ~50x to the memory floor is per-step engine overhead and not this scheme ([#2233](https://github.com/mudler/vllm.cpp/issues/2233)). q/k/v and gate/up run as separate GEMMs rather than one merged operand | +``` + +## docs/USAGE.md + +```text +| DeepSeek-V4-Flash EXL3 trellis shard 1 of 172 | `exl3-layer-000-tp4-rank0.safetensors` | 515,850,920 bytes | `0xSero/deepseek-v4-flash-0731-spark` @ `22f28d32b9b29b4352eaa380ff8c2c170b2847ab` | `2ed7ae798a794019810b027fe2609e2cf4ad78d70b49c47b2970d03a0a7aaadf` | The rank-sliced EXL3 routed-expert tower LOADS (TP4 coalesced to TP1) and its experts EXECUTE through `vt::Exl3Gemm` on a CPU queue | The CUDA arm compiles for `sm_121a` and its numeric gates PASSED on GB10 on 2026-08-28 (`had_r_128` byte-identical, `exl3_gemm` `rel_rms 5.538e-4`, GEMV tier 3c `5.160e-4`); the FUSED MoE device arm still cannot run, because it needs a device-resident tower, so the routed experts execute on a CPU queue. That run decoded ZERO tensors of THIS artifact -- it found no readable shard -- so nothing here is a claim about these weights on a device. A SYNTHETIC rank-sliced checkpoint now loads and emits logits end to end; THIS artifact still does not, because its DSA compressor and indexer tensors are stored at twice the width the host forward indexes (`compressor.wgate` `[2*head_dim, H]`) and the loader refuses them BY NAME, and because its tokenizer is not read ([#1924](https://github.com/mudler/vllm.cpp/issues/1924)) | +``` + +## docs/USAGE.md + +```text +| DeepSeek-V4-Flash EXL3 carried tower shard 1 of 5 | `carried-001.safetensors` | 4,288,630,252 bytes | `0xSero/deepseek-v4-flash-0731-spark` @ `22f28d32b9b29b4352eaa380ff8c2c170b2847ab` | `3b67ae29f1e75c2ecadfcafd3b0eecec640b06fd60b832f77e6bd3c2a8c85ccf` | The un-requantized `deepseek_v4_fp8` attention, router, shared-expert, compressor and embedding tensors, MATERIALIZED at load into the host-float tower the forward composes with — block-wise FP8 (`F8_E4M3` + `F8_E8M0` over 128x128 blocks) decoded to f32, BF16 norms and embeddings widened, I64 `tid2eid` narrowed to int32 | The DSA compressor and indexer tensors of this artifact are `2 * head_dim` / `2 * index_head_dim` wide and the loader refuses them by name (41 of its 43 layers carry a compressor); the 3,985 `mtp.*` NVFP4 draft tensors are skipped and counted, never silently dropped | +``` diff --git a/.agents/issues/QUANT-EXL3/ISSUE-LOCAL-01M2CBGG169HDX0DWKHQWCBBGN.md b/.agents/issues/QUANT-EXL3/ISSUE-LOCAL-01M2CBGG169HDX0DWKHQWCBBGN.md new file mode 100644 index 000000000..98ebe4537 --- /dev/null +++ b/.agents/issues/QUANT-EXL3/ISSUE-LOCAL-01M2CBGG169HDX0DWKHQWCBBGN.md @@ -0,0 +1,19 @@ +ID: ISSUE-LOCAL-01M2CBGG169HDX0DWKHQWCBBGN +Title: Align EXL3 user documentation with current dispatch +Row: QUANT-EXL3 +State: OPEN +Kind: docs +GitHub: - +Mirror: PENDING +Availability: FULL +Created: 2026-09-13 +Updated: 2026-09-13 +Closed: - + +## Problem + +README describes older CUDA EXL3 coverage. The EXL3 feature cell mixes current behavior with superseded attempts. Reconcile README, FEATURES and the EXL3 usage entries against source without new performance or runtime claims. + +## Resolution + +- diff --git a/.agents/specs/exl3-public-docs.md b/.agents/specs/exl3-public-docs.md new file mode 100644 index 000000000..880a55b6c --- /dev/null +++ b/.agents/specs/exl3-public-docs.md @@ -0,0 +1,138 @@ +# EXL3 public documentation + +## Now + +Documentation correction for `QUANT-EXL3`, scoped to +`ISSUE-LOCAL-01M2CBGG169HDX0DWKHQWCBBGN`. The implementation row stays `ACTIVE`. +Base: `e030f1b90`. One pull request carries this spec and the documentation. + +## Scope and inventory + +| Surface | Current gap | Source and evidence | Change | +|---|---|---|---| +| README news | The older CUDA description omits mul1 and long-prefill dispatch | `src/vt/cuda/cuda_exl3.cu:2132`, `include/vllm/model_executor/models/dense_attn_block.h:315` | Replace the old EXL3 entry with a September capability update | +| `docs/FEATURES.md` EXL3 row | Attempt history obscures current support and contains a stale HumanEval obligation | `.agents/specs/quant-exl3-shared.md`, `.agents/specs/quant-exl3-perf.md`, `docs/benchmarks/qwen38-27b-exl3-gb10.md` | Concise support summary with evidence links | +| `docs/USAGE.md` EXL3 artifact entries | Current instructions mix with old attempts | Shared dense dispatch, EXL3 loader, and the quantization specs | Retain artifact pins, hashes, sizes, working paths, and refused arms. Remove obsolete chronology | + +## Design and upstream anchors + +This is an editorial correction, not a port. The existing `QUANT-EXL3` inventory +and committed implementation spec own the upstream surface and kernel gates. +The reference is exllamav3 at the revision in `.agents/oracles/exllamav3.md`. +`quant-exl3-shared.md` records the upstream dispatch and dependency chain. +Read the local loader, dispatch, CUDA instantiations, and CPU regression test +before editing their descriptions. Do not infer full model support from a kernel. + +CUDA uses reconstruct plus cuBLASLt above 144 input rows when registered. +Other backends retain `Exl3Gemm`. Preserve the distinction between GEMM, GEMV, +and the DeepSeek-V4 fused MoE path. Do not generalize CUDA coverage to other +backends. Do not add new measured performance claims. + +Preserve removed historical text in a dated file under `.agents/completed/`. +Keep the public feature row readable and link detailed constraints to the +existing specs. Change no benchmark values or benchmark dispositions. + +## Tests and gates + +No upstream tests to port: no executable behavior changes. No GPU, model +weights, external compute, or runtime benchmarks are required or authorized. +Use source inspection to verify every changed technical statement. Retain all +artifact filenames, sizes, revisions, and hashes byte-for-byte. + +Focused CPU gate: + +```sh +python3 scripts/check-readme-structure.py +python3 scripts/check-supported-models.py +python3 scripts/check-quickstart-recipes.py +python3 scripts/check-benchmark-index.py +python3 scripts/check-agent-record.py +python3 -m unittest discover -s tests/scripts -p 'test_check_readme_structure.py' +git diff --check +``` + +Run `scripts/agent-preflight.sh --quiet` and classify any failure against the +unchanged base. An independent reviewer checks the immutable implementation +commit and reruns the focused gate. Mutations apply only to executable test +claims. No new tests that merely copy the prose are needed for this edit. +The operator reruns the focused gate before the fork push. + +## Work breakdown and authority + +1. Operator commits this spec and the canonical local issue. +2. A fresh implementer edits only the three public documents above, the dated + archive, and this spec's outcome. It records source anchors and gate output. +3. A fresh reviewer checks the immutable head without repairing findings. +4. Operator verifies and opens a pull request from the maintenance fork to + `mudler/vllm.cpp:main`. Merge is not authorized by the user request. + +## Risks and stop conditions + +Do not promote a single completed benchmark round to an accepted headline. +HumanEval-style benchmark data is not a correctness gate. Preserve those limits. +Stop on an ambiguous source claim rather than guessing. Do not modify source, +tests, checkers, other rows, or unrelated documentation to repair baseline gates. + +## Outcome: 13 September 2026 + +Implemented the three public-document corrections without changing model +support states or benchmark values. Removed prose is preserved verbatim in +[the dated archive](../completed/exl3-public-docs-20260913.md). The five edited +artifact rows retain their first five cells byte-for-byte, including filenames, +sizes, revisions, and hashes. The CUDA news replaces the older August entry. + +Source verification at implementation base `0cf30baae`: + +| Statement | Inspected source or existing evidence | +|---|---| +| Per-tensor width and codebook selection | `include/vllm/model_executor/models/dense_weight_loaders.h:700` derives bits from trellis shape; `:750` reads `mcg` and `mul1` presence | +| Shared dense method and long-prefill dispatch | `include/vllm/model_executor/layers/quantization/exl3.h:70` delegates to `dense_attn::Exl3MatmulD`; `include/vllm/model_executor/models/dense_attn_block.h:315` requires both M > 144 and a registered reconstruct op | +| CUDA GEMM coverage | `src/vt/cuda/cuda_exl3.cu:2132` admits exactly `(3,0)`, `(3,1)`, `(6,0)`, `(3,2)`, `(4,2)`, `(5,2)`, and `(6,2)`; reconstruct dispatch at `:2707` repeats those pairs | +| GEMV is a separate capability | `src/vt/cuda/cuda_exl3.cu:2284` admits `(3,1)`, `(3,2)`, and `(4,2)`; `:2314` declines uninstantiated arms; `tests/vt/test_exl3_gemv.cpp:161` covers artifact shapes and occupancy constraints | +| Non-CUDA long-prefill behavior | `tests/vt/test_exl3_matmul_dispatch.cpp:68` exercises M=145 through the shared seam on CPU, with no reconstruct registration. Inspected here, not compiled or executed | +| Draft uses packed target head | `src/vllm/model_executor/models/qwen3_dflash.cpp:105` selects the EXL3 head before the other formats | +| DeepSeek-V4 doubled widths no longer fail at load | `src/vllm/model_executor/models/deepseek_v4_weights.cpp:1145` derives compressor width from the ratio; `:1195` loads the indexer compressor at twice its head dimension | +| DeepSeek-V4 tokenizer caveat is obsolete | `tests/vllm/test_tokenizer_parity_deepseek_v3.cpp:136` compares encoding with HF goldens; `:144` checks round trips. `model-dsv4-exl3.md:53` records the real artifact's tokenizer parity evidence | +| DeepSeek-V4 MTP exclusion remains | `src/vllm/model_executor/models/deepseek_v4_weights.cpp:1306` explicitly skips and counts `mtp.*` tensors | +| Generation and performance limits | `docs/benchmarks/qwen38-27b-exl3-gb10.md:3` records generation; `:27` excludes correctness claims for sampled HumanEval legs. `quant-exl3-perf.md:45` records GEMV occupancy limits; `backend-rocm-exl3.md` owns gfx1151 generation evidence | + +Focused verification uses Python from `/tmp/vllm-docs-tools/usr/bin`, with +`LD_LIBRARY_PATH=/tmp/vllm-docs-tools/usr/lib`: + +| Command | Result | +|---|---| +| `python3 scripts/check-readme-structure.py` | PASS, exit 0 | +| `python3 scripts/check-supported-models.py` | PASS, exit 0; 44 registered architectures | +| `python3 scripts/check-quickstart-recipes.py` | PASS, exit 0 | +| `python3 scripts/check-benchmark-index.py` | PASS, exit 0 | +| `python3 scripts/check-agent-record.py` | PASS, exit 0 after correcting the new claim's missing lifecycle annotation; initial exit 1 | +| `python3 -m unittest discover -s tests/scripts -p 'test_check_readme_structure.py'` | PASS, exit 0; 19 tests | +| `git diff --check` | PASS, exit 0 | + +A direct comparison with `git show 0cf30baae:docs/USAGE.md` confirms the five +artifact rows retain their first five cells. A second comparison confirms every +removed FEATURES and USAGE line occurs verbatim in the archive. Both pass. + +No executable behavior or test guarantee changes, so red-first and negative +mutation checks are not applicable. No GPU or runtime benchmark was run. The +operator owns the unchanged-base comparison and full `agent-preflight.sh` +result; its broad run was still in progress at this implementation handoff. +Missing compiler, CMake, and readelf are reported environment limitations, +not passing gates. This helper does not modify unrelated code to repair them. + +### Review corrections: 13 September 2026 + +The MTP bit width applies to its trellis modules, not every companion tensor. +`quant-exl3-mul1.md:246-248` records eight 4-bit MTP modules. +`dense_weight_loaders.h:722-723` requires F16 sign vectors, and `:765-767` +requires an I32 `mul1` marker. The usage description now names the trellis +modules explicitly. + +The ROCm evidence includes a subsequent successful BF16 control at +`backend-rocm-exl3.md:81-104`. The usage limitation now names missing AMD clock +attribution and discrete-GPU validation. No benchmark values or artifact pins +changed. Both corrections address the existing documentation issue. + +All seven focused gates passed again after these corrections, including all +19 README checker tests. The first five cells of both artifact rows remain +byte-identical to repair base `752b9628b`. No GPU work was run. diff --git a/README.md b/README.md index 3f80e29e8..e3912c18b 100644 --- a/README.md +++ b/README.md @@ -43,9 +43,10 @@ - **2026-09** **EXL3 gains a native ROCm path.** Llama-3.2-1B-Instruct EXL3 generated on gfx1151 with zero CPU fallbacks. Discrete AMD validation and competitive performance remain unmeasured. See the [ROCm build guide](docs/BUILD.md#rocm-build-amd-gpus). -- **2026-08** **EXL3 checkpoints now generate on CPU and CUDA.** A stock - Llama-3.2-1B-Instruct EXL3 checkpoint loads through the shared dense model path and emits text. - The current CUDA path supports its 3-bit body and 6-bit output head. No speed claim is available. +- **2026-09** **CUDA EXL3 now supports Qwen3.8-27B and its DFlash2 draft.** + The CUDA path handles their `mul1` codebooks and uses reconstruction plus cuBLASLt for long + prefills. See [supported artifacts and limits](docs/USAGE.md#checkpoint-registry) and + [benchmark conditions](docs/benchmarks/qwen38-27b-exl3-gb10.md). - **2026-08** **GLM-5.3-Flash now generates on CPU from a 101.25 GiB GGUF.** The shipped `UD-Q2_K_XL` artifact emits coherent text while keeping IQ2_XS and IQ4_XS blocks compressed. Both formats also have CUDA keep-quant kernels, but this model's CUDA forward and every speed diff --git a/docs/FEATURES.md b/docs/FEATURES.md index ae3f28dc1..366418134 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -112,7 +112,7 @@ by #2794 (goldens predate the pin) and #2817 (the advance). | GGUF gather tables kept QUANTIZED (one row dequantized per gathered token) | ✅ CPU, `qwen35`/`qwen35moe`/`qwen3next`/`qwen4exp`. `vt::Embedding` takes a block-quantized table, decoding one row per id — a port of llama.cpp's `ggml_compute_forward_get_rows_q`. Without it a 51.2 G-parameter n-gram table expands from 28.8 GB of IQ4_NL to 102.4 GB of bf16, which no device here has (#1989). This is a residency DEFAULT CHANGE on already-shipped GGUF models, not only a new arm: an existing `qwen35` file with a quantized `token_embd` now keeps it compressed on CPU where it used to expand. Tokens do not move (every GGUF-path gather writes a bf16 output and the bf16 round is idempotent over the old expand-then-widen), so the change is memory-only today. `deepseek4` and `laguna` are NOT reached: both consume `token_embd` as a flat host f32 array, so their loaders narrow the policy for that tensor and keep expanding it. A gather's admission is the ROW DECODER and not the `vec_dot`, so IQ2_XS and IQ4_XS joined this set the moment their decoders landed (#2240), a wave before they could keep on the GEMM arm; Q8_K is now the only encoding that still tells the two admission rules apart. **CUDA joined it** (KGATHER, [spec](../.agents/specs/cuda-quant-gather.md)). `EmbeddingKernelCuda` decodes a block row across all 19 encodings `vt::cpu::BlockToFloat` decodes (IQ3_S joined with [#2510](https://github.com/mudler/vllm.cpp/issues/2510)), so a gather table stays block-resident ON THE CARD. The gate is a REGISTRY QUERY and not a device list: `vt::Embedding` routes a block table to `OpId::kEmbeddingQuant` and `DeviceQuantGatherSupported` is `OpRegistered(kEmbeddingQuant, dev)`, so a backend advertises the capability by registering the kernel and a device with no decoder refuses BY NAME. **MEASURED ON A GPU**, nvcc 13.0.88, 2026-08-31, and the two boxes measured DIFFERENT things: **thor:gpu0 (Jetson Thor, sm_110)** ran the branch HEAD through the production registration — 7 of 7 cases at 250 assertions, every encoding bit-exact against the CPU arm in f32 and bf16 out with i32 and i64 ids, against a pre-arm RED leg that failed by name at 32 assertions, and deleting the single production registration reds 5 of 7 cases, so the arm is reached and not merely present. **dgx:gpu0 (NVIDIA GB10, sm_121a)** then ran the SAME head leg green — 7 of 7 at 250, M5 red, every verdict matching thor — so the arm is proven on BOTH architectures and [#2393](https://github.com/mudler/vllm.cpp/issues/2393) is closed. This is what makes a GPU arm of `qwen4exp` possible at all — its n-gram table is 26.822 GiB of IQ4_NL against 95.368 GiB expanded, on a box with ~119.6 GiB for everything — and it is a shape llama.cpp does NOT have: its CUDA `get_rows` dispatches the legacy quants only and aborts on every K-quant and every IQ type, so the llama.cpp column overstates the denominator for the encodings that matter here. Like the CPU arm this is a residency DEFAULT CHANGE on already-shipped GGUF models, and it reaches more than `qwen4exp`: a `token_embd` in a block encoding now keeps its blocks on CUDA where it used to expand, `qwen35` GGUFs included. **For a bf16 output the values do not move** (both routes decode with the same scalar expressions and round once); **for an f32 output the new path is strictly MORE precise than expand-then-widen**, which is a behaviour change on already-shipped models and not merely a memory one. **ROCm now registers a native HIP gather for all 19 admitted encodings** ([spec](../.agents/specs/rocm-quant-gather.md), [#3093](https://github.com/mudler/vllm.cpp/issues/3093)). The gfx1100 public gate passes each codec against its dense control, and matching operation traces retain compressed tables with direct bf16 output. Original primary fixtures and pinned model qualification remain pending. **Owed: the gather arms of METAL, VULKAN and TENSTORRENT ([#2394](https://github.com/mudler/vllm.cpp/issues/2394))**, whose gather kernels still require floating tables. No throughput number is claimed | ☐ | ☐ | ✅ `get_rows` for ~20 types, CPU and CUDA | | GGUF F16 weights kept resident as F16 | ✅ CPU `VT_GGUF_KEEP_F16` defaults on and computes directly on F16. The measured CPU trade is 1.05 GiB less peak RSS for ~9% prefill and ~1.4% decode, with identical tokens. ◐ ROCm Qwen3.5 dense retains eligible projections and embedding tables by default, then rounds their values to the resolved BF16 model dtype before arithmetic. Ordinary F16 GEMM and gather checks pass on gfx1100. The complete model gate remains failing, and no ROCm memory or speed improvement is accepted ([spec](../.agents/specs/rocm-f16-weights.md), [#3092](https://github.com/mudler/vllm.cpp/issues/3092)). `VT_GGUF_KEEP_F16=0` restores expansion. Stacked experts and unwired loaders retain expansion | ☐ | ☐ | ✅ `ggml_vec_dot_f16` | | GGUF is a TWO-engine comparison at these pins (#979) | ✅ text-only `qwen35`, no `clip` projector (#821) | ☐ REMOVED from the tree in `6635279d8`, now an unpinned out-of-tree `vllm-gguf-plugin` | ☐ full stack present, `qwen3_5` unreachable behind FOUR blockers, and the load path has NO completeness guard so a clean-looking load proves nothing | ✅ native, `LLM_ARCH_QWEN35` | -| EXL3 trellis (exllamav3; codebooks 0 (3INST), 1 (MCG) and 2 (`mul1`), Hadamard-128 + sign vectors, NO scales) | ◐ **A stock EXL3 checkpoint GENERATES**: `turboderp/Llama-3.2-1B-Instruct-exl3` @ 3.0bpw loads through the shared dense container and emits coherent text from `vllm-cli` on a CUDA queue and on a CPU queue, which reaches `LlamaForCausalLM` and Qwen3-dense together. The scheme sits on vLLM's own `LinearMethodBase` seam, so it is no longer a DeepSeek-V4-private arm. `bits` and the codebook are both read PER TENSOR — the published 3.0bpw Llama has a 3-bit body and a 6-bit head, and ships no `mcg` marker, which means codebook 0 and not MCG. **The DeepSeek-V4 arm is separate and unchanged**: its rank-sliced SparkInfer artifact loads and executes end to end on a synthetic checkpoint. The rank-sliced `0xSero/deepseek-v4-flash-0731-spark` routed-expert tower coalesces TP4->TP1 at load, the `carried-*` half (block-wise FP8 + BF16 + F32) is dequantized into the host-float tower the forward composes with, and `DeepseekV4Model::Forward` runs the whole model, dispatching one `vt::Exl3MoeMlp` per MoE layer over the routed experts (`bits == 3`, `mcg` codebook). **The REAL artifact still does not run**: its DSA compressor and indexer tensors are twice the width the host forward indexes and the loader refuses them by name, and its tokenizer is not read (#1924). **The device half is now PARTLY VERIFIED** (GB10 `sm_121a`, 2026-08-28): `had_r_128` is BYTE-IDENTICAL CUDA-vs-CPU, `exl3_gemm` matches the f64 reference at `rel_rms 5.538e-4` against a `1.0e-3` bound, and the `m<=8` GEMV meets tier 3c at `5.160e-4` against `6.0e-3`. What is STILL unverified on a device is the FUSED MoE arm, which cannot run on this code at all: it needs a device-resident tower and `CudaBackend::DeviceMemoryIsHostAddressable()` is false by design, so the routed-expert path executes on a CPU queue today. No speed number is claimed on any axis. **Codebook 2 (`mul1`) and the 4- and 5-bit widths are now implemented** so that `Mia-AiLab/Qwen3.8-27B-EXL3-3.5bpw` -- **409** trellis modules, 137 at 3 bpw, 270 at 4, one at 5 and one at 6, every one of them `mul1`-marked -- is no longer refused by name ([#2495](https://github.com/mudler/vllm.cpp/issues/2495)). The count read 272 until 2026-09-02 and was an UNDERCOUNT that omitted the 137 bits-3 modules, which is why `(3, 2)` was missing from the CUDA instantiation list and every MLP projection of that target refused on a CUDA queue. **`(3, 2)` is now instantiated and gated on `dgx:gpu0`** (GB10, driver `580.173.02`), agreeing with the CPU decoder at `rel_rms 3.02544e-07` against a `1.0e-3` bound ([#2574](https://github.com/mudler/vllm.cpp/issues/2574)). cb 2 is a DIFFERENT decode and not a third multiplier: it sums the four bytes of the product into an fp16 bit pattern and maps it with a fused fp16 affine, and the host arm is gated against hand-computed upstream values rather than against itself. **A ROCm arm now exists too** ([#2433](https://github.com/mudler/vllm.cpp/issues/2433), [spec](../.agents/specs/backend-rocm-exl3.md)): `Exl3Gemm` on `gfx1151` is a transcription of the PORTABLE CPU reference rather than a port of the CUDA kernel -- 90 KiB of requested shared memory does not fit a 64 KiB AMD LDS, and `m16n8k16`, `ldmatrix` and `cp.async` have no AMD form -- and because it uses no matrix cores it is gated at BYTE equality with the CPU arm on all three codebooks, not at a tolerance. It covers every width the host arm decodes, so the CUDA instantiation list below does not bound it. On the CUDA device arm `(3,0)`, `(3,1)`, `(6,0)`, `(3,2)`, `(4,2)`, `(5,2)` and `(6,2)` are instantiated -- `(3,2)` is the width the undercount cost, added and device-gated by [#2574](https://github.com/mudler/vllm.cpp/issues/2574); every other width and codebook and the device-resident tower are owed ([spec](../.agents/specs/quant-exl3-mul1.md)). **The CUDA GEMM now has a dual-path dispatch mirroring exllamav3** ([#3124](https://github.com/mudler/vllm.cpp/issues/3124), [spec](../.agents/specs/quant-exl3-shared.md)): M <= 144 routes to the cooperative `exl3_gemm` kernel; M > 144 routes to `Exl3ReconstructGemm`, which dequantizes trellis weights to fp16 on-device and runs cuBLASLt fp16 GEMM. Only CUDA registers the reconstruct kernel, so CPU, ROCm and Vulkan keep `exl3_gemm` at every M, as they did before the dispatch existed. The fused path (M >= 1024, both dims 128-divisible) folds the input/output Hadamards into the reconstruct kernel; the unfused path applies them separately. N > 32768 slices the output dimension. All 7 (bits, codebook) arms are gated against an f64 reference at `rel_rms <= 1e-3` on `dgx:gpu0`. **The `m<=8` GEMV is a SEPARATE arm set from that list and a separate row** ([`QUANT-EXL3-PERF`](../.agents/specs/quant-exl3-perf.md), [#2570](https://github.com/mudler/vllm.cpp/issues/2570)): upstream instantiates `(4,0) (4,1) (4,2) (2,1) (2,2) (3,1) (3,2)` and this tree now carries `(3,1)`, `(3,2)` and `(4,2)`, so **407 of that checkpoint's 409 trellis modules** reach the small-m fast path where before NO module of it could -- the only arm was `(3,1)` and the artifact contains zero `(3,1)` tensors. `(3,2)` was one template argument; `(4,2)`, the 270-module arm, was a KERNEL port: a per-width `LSTRIDE`, two 24-lane guards that exist only because a 3-bit tile is 24 uint32 rather than 32, and a separate `dq8_regs_4bits` window read. `(4,0)`, `(4,1)` and the whole 2 bpw arm stay OWED -- no artifact in this tree carries a 4-bit tensor at either of the first two codebooks, and nothing builds them, so that the kernel would compile for them is an inference and not a measurement. **Instantiating an arm is necessary and not sufficient**: `Exl3GemvSelectConfig` is upstream's envelope verbatim and DECLINES, and on Blackwell EVERY shape of this checkpoint, at both widths, is admitted only where `size_n / 32 <= narrow_coresident`, an occupancy query. The wide config is not an escape from it: its band needs `size_k <= 4096` and the smallest 4-bit `k` here is 5120. The thresholds are `>= 544` and `>= 160` at the two bits-3 shapes and 32, 160, 192, 320, 384 and 544 across the six bits-4 ones, all pinned in `tests/vt/test_exl3_gemv.cpp` from both sides. **The throughput effect is now MEASURED on GB10, and it is a NULL** ([#2570](https://github.com/mudler/vllm.cpp/issues/2570)): three interleaved rounds on one binary read 17.20/17.18/17.09 tok/s with the arm off against 17.22/17.12/17.09 with it on, a 1.3% spread and no separation. `nsys --cuda-graph-trace=node` shows why, and it is a DECLINE and not an ineffective kernel: `exl3_gemv_kernel` appears 0 times at the default and 4 times only when forced. GB10 reports `SM_COUNT=48 MAX_THREADS_PER_SM=1536`, so the narrow config's 512-thread blocks ceiling `narrow_coresident` at 144, and only shapes needing `<= 144` are admitted -- 34 of the checkpoint's 409 modules, 0.75% of its trellis bytes. The arms are correct, upstream-faithful and device-gated; on THIS part the envelope declines them. `m == 8` is unmeasured, because the draft legs did not run. **The FUSED MoE mgemm is a SEPARATE arm set from that list**, and upstream's own bound on it is narrower: `exl3_moe.cu:184` admits `mcg` and `mul1` and refuses 3INST, so codebook 0 is not a MoE arm upstream either. It now instantiates bits 3, 4, 5 and 6 at codebook 1 where it carried bits 3 alone ([#2756](https://github.com/mudler/vllm.cpp/issues/2756)); before that a mcg expert tower at any other width could not run on a CUDA queue by ANY path, because the fused arm is default-ON and its refusal is not caught, and the `VT_DSV4_EXL3_FUSED_MOE=0` rollback lands on a GEMM arm set with no `(4,1)`, `(5,1)` or `(6,1)` either. Codebook 2 on the fused arm is owed as a LOADER slice and not a kernel one: five sites pin the codebook to 1 before the kernel is chosen, so a cb-2 instantiation added alone would be dead code. **Its GDN linear-attention tower now loads and runs too** ([#2495](https://github.com/mudler/vllm.cpp/issues/2495) item 4, [spec](../.agents/specs/model-qwen35-gdn-exl3.md)): 48 of that model's 64 layers are `linear_attention`, and the three projections of each were refused by name because nothing in the GDN forward consumed an `Exl3Weight`. The arm holds THREE trellises rather than one merged `in_proj_qkvz` owner, because the artifact ships `in_proj_qkv` and `in_proj_z` as two independently quantized tensors whose sign vectors are fitted per projection, so the bf16 merge has no trellis analogue. It also accepts the F16 `in_proj_a`/`in_proj_b` the artifact stores beside its BF16 `conv1d`, `norm`, `A_log` and `dt_bias`. **The DFlash2 DRAFT loader now has an EXL3 arm too** ([#2495](https://github.com/mudler/vllm.cpp/issues/2495) item 7, [spec](../.agents/specs/model-dflash2-exl3.md)): `Mia-AiLab/Qwen3.8-27B-DFlash2-EXL3-5.0bpw` ships all 36 of its modules at bits 5 with a `mul1` marker, and `LoadQwen3DFlash` died on `fc.weight`, which an EXL3 module does not ship. The arm holds SEVEN trellises per layer rather than the merged `qkv_proj` and `gate_up_proj` owners the bf16 reader builds, for the same reason the GDN tower holds three. It also accepts the F16 `candidate_selector.hidden_projection` and `*_conv.kernel_projection` the repack left unquantized beside their BF16 `base_kernel` and codebooks — a second refusal that was invisible until the first was gone. **Its `mtp.*` draft head loads and runs too, and all THREE readers of the target's trellis head now compute with it packed** ([#2495](https://github.com/mudler/vllm.cpp/issues/2495) items 5 and 6, [spec](../.agents/specs/model-qwen35-exl3-head.md)): the MTP loader refused nine quantized `mtp.*` tensors as "expected BF16", and the MTP draft's `ComputeLogits` plus the DFlash/DSpark draft's SHARED head read knew only bf16 and NVFP4. The head is COMPUTED WITH rather than widened, which is what upstream does without a branch because its head is an `nn.Module` and `_apply_head` calls `lm_head.quant_method.apply`; a dequantized copy of the real 248320x5120 head would be 2.543 GB against 0.953 GB packed. The shared read also used to fall off the end of its bf16 loop and return SILENTLY with every owner empty when no arm matched, which is exactly what this target does -- it ships no `lm_head.weight` at all -- and it now refuses by name ([#2569](https://github.com/mudler/vllm.cpp/issues/2569)). **The checkpoint's real inventory is 409 quantized modules, not the 272 this row previously recorded**, measured from both downloaded shards' headers: bits 3 x137, 4 x270, 5 x1, 6 x1, every one `mul1`, and no tensor in the file has `scale` in its name. `(bits 3, codebook 2)` is now instantiated on the CUDA GEMM and gated on a device, so those 137 modules no longer refuse by name ([#2574](https://github.com/mudler/vllm.cpp/issues/2574), owned by `QUANT-EXL3-MUL1`, not by this row). **THE 27B CHECKPOINT NOW RUNS ON CUDA AND GENERATES.** On `dgx:gpu0` (GB10, driver 580.173.02, `sm_121a`) the published `Mia-AiLab/Qwen3.8-27B-EXL3-3.5bpw` loads and emits ` Paris. The capital of Germany is Berlin. The capital of Italy is` at `rc=0`, with both shard sha256 values recomputed on the device. Warm decode is **16.7 tok/s** target-only and **48.7 tok/s** with its DFlash2 EXL3 draft at k=7, interleaved on one binary, and the two arms are **token-identical** ([`qwen38-27b-exl3-gb10`](benchmarks/qwen38-27b-exl3-gb10.md)). **That is NOT a reproduction of the upstream README's 47.5 tok/s**, and must not be quoted beside it: theirs is HumanEval-style at T=0.6 with acceptance 4.43, this is greedy T=0 on a prompt that continues into a list of capitals, the context is 8192 rather than 262144 with no NVFP4 KV cache, and the shipped paged draft route had to be disabled (#2274). A matched HumanEval-style run at T=0.6 with acceptance reported is owed | ☐ no EXL3 at the parity pin | ☐ | ☐ | +| EXL3 trellis (exllamav3; codebooks 0, 1, and 2) | ◐ Llama-3.2-1B generates on CPU, CUDA, and ROCm gfx1151. Qwen3.8-27B and its EXL3 DFlash2 draft generate on CUDA. The shared dense path reads each tensor’s bit width and codebook. CUDA reconstructs weights for cuBLASLt above 144 input rows; CPU, ROCm, and Vulkan retain `Exl3Gemm`. CUDA GEMM, GEMV, and fused MoE have different coverage. The DeepSeek-V4 real-artifact gate remains open. See [artifacts and limits](USAGE.md#checkpoint-registry), [dispatch coverage](../.agents/specs/quant-exl3-shared.md), [GEMV limits](../.agents/specs/quant-exl3-perf.md), and [measured workloads](benchmarks/qwen38-27b-exl3-gb10.md). The sampled HumanEval comparison does not establish correctness. | ☐ no EXL3 at the parity pin | ☐ | ☐ | | AWQ | ◐ CPU dequant | ✅ | ✅ | ☐ | | GPTQ | ◐ CPU dequant | ✅ | ✅ | ☐ | | MXFP4 compressed-tensors | ◐ W4A16 Marlin, mem 2.63x less. gate_up FUSION + decode-graph default-ON; #44 3/3, 32B 6/6. **`VT_MARLIN_DENSE` DEFAULT-ON** (`KERNEL-MARLIN-DENSE-EXEC`): dense marlin 48-CTA, byte-faithful, beats MoE (c8 0.969) | ✅ | ✅ | ☐ | diff --git a/docs/USAGE.md b/docs/USAGE.md index b72c953d5..5a1a6c2f8 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -978,10 +978,10 @@ repository in this project's history. |---|---|---|---|---|---|---| | Qwen3.5-0.8B text, ROCm retained F16 validation | `Qwen3.5-0.8B-F16.gguf` | 1,557,662,528 bytes; 335 tensors | Locally converted from `Qwen/Qwen3.5-0.8B` @ `2fc06364715b967f1860aea9cf38778875588b17` using llama.cpp `10bf611e533d81f739128304991c5e133c6aebd8` | `758b5299b027120c3608c43777a89257724d46e26eefac2a8aaad512be15b53f`, measured from the generated file | Default public loading and generation reach retained F16 projections and embedding on gfx1100. The [token gate remains failing](../.agents/specs/rocm-f16-weights.md), so this is a validation artifact with no accepted performance result | Text-only GGUF omits the vision projector. F16 model activations and outputs are not enabled. Stacked expert F16 retention remains unimplemented. Other registry loaders retain expansion | | DSpark for Qwen3.8-27B | `model.safetensors` | 2,718,576,122 bytes | `RadixArk/Qwen3.8-27B-DSpark` @ `85ef153be924f17ce4bf62726954eeaa4a73e854` | n/a (non-quantized) | Qwen3 DSpark routing | Token-exact decode gate is pending | -| Qwen3.8-27B EXL3 3.5bpw (the #2495 benchmark target; RUNS and GENERATES on CUDA) | `model-0000{1,2}-of-00002.safetensors` | 15,338,408,461 bytes total (index `total_size` 15,338,106,948 plus the two safetensors headers); 2426 tensors | `Mia-AiLab/Qwen3.8-27B-EXL3-3.5bpw` @ `19441ac874c4018295da848e250f23511361cda4` | `7b77214fe58ff15fed0b4af55e3cd92f38842b8711886d68954e8071ff8270c6` (shard 1), `411c83bb1070b27f3d670fc93e38dca0f17eb66429f64b5706901b12613188b2` (shard 2) — both **recomputed on `dgx:gpu0`** and matching the download host | `Qwen3_5ForConditionalGeneration`, 64 layers as 48 `linear_attention` + 16 `full_attention`. **409 trellis modules, every one codebook 2 (`mul1`), zero `mcg`**: bits 3 x137 (`mlp.{gate,up,down}_proj` of 46 layers), bits 4 x270 (all 48 GDN `linear_attn` x3, 55 more `mlp.*`, the 16 `self_attn.*`), bits 5 x1, bits 6 x1 (`lm_head`). It also ships its own quantized MTP head, 39 tensors under `mtp.*`, uniformly bits 4. GENERATES on GB10 greedy at `rc=0`; 16.7 tok/s target-only and 48.7 tok/s with its DFlash2 draft ([`qwen38-27b-exl3-gb10`](benchmarks/qwen38-27b-exl3-gb10.md)) | Its vision tower (`language_model_only: false`, depth 27) is not exercised. `quantization_config` says `out_scales: "always"` but the artifact ships **no scale tensors** — every quantized module is exactly `{trellis, suh, svh, mul1}` and the setting is folded into `svh`. The `m <= 8` GEMV had an arm for **none** of its 409 modules until 2026-09-02 -- its only arm was `(3,1)` and this artifact has zero `(3,1)` tensors. `(3,2)` is now instantiated, which is the 137 bits-3 modules, and `(4,2)` too, which is the other 270 ([#2570](https://github.com/mudler/vllm.cpp/issues/2570), [`QUANT-EXL3-PERF`](../.agents/specs/quant-exl3-perf.md)); 407 of the 409 now have an arm, and the 5- and 6-bit pair have none upstream either. **Whether the arm is TAKEN is a separate question from whether it exists**: `Exl3GemvSelectConfig` is upstream's envelope verbatim and declines on Blackwell below an occupancy threshold -- `narrow_coresident >= 544` at k 5120 n 17408 and `>= 160` at k 17408 n 5120, and 32, 160, 192, 320, 384 and 544 across the six bits-4 shapes -- and that HAS now been measured on GB10: `SM_COUNT=48 MAX_THREADS_PER_SM=1536` ceilings `narrow_coresident` at 144, so 34 of the 409 modules are admitted and the decode A/B is a null (17.16 vs 17.14 tok/s mean over three interleaved rounds, 1.3% spread), with `nsys` showing zero `exl3_gemv_kernel` launches at the default. The wide config does not escape it here, because its band needs `size_k <= 4096` and the smallest 4-bit `k` in this artifact is 5120. No throughput change is claimed here | +| Qwen3.8-27B EXL3 3.5bpw (the #2495 benchmark target; RUNS and GENERATES on CUDA) | `model-0000{1,2}-of-00002.safetensors` | 15,338,408,461 bytes total (index `total_size` 15,338,106,948 plus the two safetensors headers); 2426 tensors | `Mia-AiLab/Qwen3.8-27B-EXL3-3.5bpw` @ `19441ac874c4018295da848e250f23511361cda4` | `7b77214fe58ff15fed0b4af55e3cd92f38842b8711886d68954e8071ff8270c6` (shard 1), `411c83bb1070b27f3d670fc93e38dca0f17eb66429f64b5706901b12613188b2` (shard 2) — both **recomputed on `dgx:gpu0`** and matching the download host | `Qwen3_5ForConditionalGeneration`: 48 linear-attention and 16 full-attention layers. All 409 trellis modules use codebook 2 (`mul1`): 137 at 3 bits, 270 at 4, and one each at 5 and 6. The MTP trellis modules use 4-bit weights. Target-only and DFlash2-assisted generation run on GB10; see the [benchmark record](benchmarks/qwen38-27b-exl3-gb10.md). | The vision tower is untested. No scale tensors ship; output scaling is folded into `svh`. CUDA GEMM covers all four widths. GEMV has 3- and 4-bit arms, but its occupancy rules restrict dispatch on GB10; the recorded decode comparison shows no throughput gain. See [GEMV coverage and measurements](../.agents/specs/quant-exl3-perf.md). The HumanEval comparison is a sampled performance workload, not a correctness gate. | | Qwen3.8-27B ModelOpt NVFP4, the SGLang comparator's own target | `model-0000{1,2,3}-of-00003.safetensors` | 9,965,652,544 + 9,985,757,064 + 1,970,287,672 bytes; 2194 tensors | `RadixArk/Qwen3.8-27B-NVFP4` @ `554ebba9b5f1b79dc11246341960360e6ef05ef4` | `fbcdb5ba1cdda462b5f38592d071e772c4d398afea61a0aa9188b32d1a239a79` (shard 1), `db6146a5464fb0a891181b93c81593f0ca65c602eb14120a1c2b1b09bca11f85` (shard 2) and `d3cfb92742e30c8b46564665791dbe0a86ed64cfc02b1275081530793c0c9581` (shard 3), each **computed from the local bytes** and each matching the publisher's LFS object hash. Shard 2 was fetched as five disjoint HTTP ranges and reassembled after a single stream stalled, so its hash is the check that the reassembly is the file rather than a plausible-looking one. Leg B of `/workspace/nvfp4-sota/job1.sh` recomputes all three again on the device | ModelOpt `MIXED_PRECISION`: **208 per-tensor static FP8** (`self_attn.{q,k,v,o}_proj`, `linear_attn.{in_proj_qkv,in_proj_z,out_proj}`) and **193 NVFP4 `group_size` 16** (every `mlp.{gate,up,down}_proj` plus `lm_head`), all 401 carrying an `input_scale`. Both whole-checkpoint refusal gates answer empty on this `config.json` and its 2194 shipped tensor names, and shard 1 agrees with the committed `r0b0tlab` fixture on the dtype and shape of all 970 tensors they share. **The device load is leg E of that job and has not run**, so no arm here is claimed to execute | Its 193 NVFP4 modules are declared `NVFP4`, which is W4A4, and this build takes the **W4A16 weight-only** arm against them, because routing is by tensor name and `VT_MODELOPT_W4A4` defaults to `0`. Since `QUANT-QWEN38-27B-NVFP4-ARM` W7 the load SAYS so: one stderr notice naming both algorithms, the 193 modules, the 193 `input_scale` divisors it drops and the knob. The divergence is unpaid, not silent ([#2760](https://github.com/mudler/vllm.cpp/issues/2760)). The declared `kv_cache_scheme` FP8 is unread and bf16 KV is used, which is what the comparator also runs. The 15 bf16 `mtp.*` tensors are present and MTP execution is owed. The vision tower (`language_model_only: false`, depth 27) is not exercised | | Qwen3.8-27B DFlash2 drafter, ModelOpt NVFP4 (the SGLang comparator's own drafter) | `model.safetensors` | 1,550,153,248 bytes; 186 tensors | `maurienne-ai/Qwen3.8-27B-DFlash2-NVFP4-RTNcal` @ `bd7a934213c47a9e7ef69eef36bb3325f47fd1f1` | `2228b9b22e93a88d84556419c879448ab6c490ae65c4c0b166f4962190ddbf26`, **computed from the local bytes** and matching the publisher's LFS object hash | **None.** Its `dflash_config` is byte-identical to the EXL3 DFlash2 drafter this tree does load (`block_size` 8, taps `[5,19,33,47,61]`, `selector_rank` 256, `selector_top_k` 16, `mask_token_id` 248070), and its five `layer_types` are all `sliding_attention` at `sliding_window` 2048 | **LOADS as of [#2758](https://github.com/mudler/vllm.cpp/issues/2758); NOT YET RUN on a device.** 35 of its Linears are ModelOpt NVFP4 `group_size` 16 in the ModelOpt spelling (`.weight` U8 + `.weight_scale` F8_E4M3 + `.weight_scale_2` + `.input_scale`), and the loader now reads the DRAFT's own `quantization_config` as upstream does (`get_draft_quant_config`, `vllm/model_executor/models/utils.py:929-948`) and takes the packed arm on all seven projections of each of the five layers. Its 12 `exclude_modules` -- `fc`, `candidate_selector.hidden_projection` and both conv `kernel_projection`s per layer -- are the exact four modules this engine has NO packed owner for, and a checkpoint that quantized any of them is refused BY NAME rather than by a dtype. Before that change the EXL3 probe on `fc` answered false, the BF16 reader was chosen, and the load died inside layer 0 with `qwen3_dflash: expected BF16 for layers.0.self_attn.q_proj.weight`. It declares `quant_algo: "NVFP4"` -- W4A4 -- and this build runs W4A16, which the load now SAYS ([#2760](https://github.com/mudler/vllm.cpp/issues/2760)). The verbatim device string is still leg F of that job; no GPU has loaded it | -| Qwen3.8-27B DFlash2 draft, EXL3 5.0bpw (the drafter for the row above) | `model.safetensors` | 1,470,916,078 bytes; 189 tensors | `Mia-AiLab/Qwen3.8-27B-DFlash2-EXL3-5.0bpw` @ `4f0436269bca761b071f05319e8e04a87cc633f9` | `6b2e3afc694a343b7f3f0edfe5925e460762fc9ede4699165b577ca0733c8e56` | `DFlash2DraftModel`, 5 sliding-attention layers, `block_size 8`, taps at target layers `[5,19,33,47,61]`. **36 trellis modules, all bits 5 codebook 2** — uniform, NOT the "module-adaptive" the upstream README claims. It owns no `lm_head` and no `embed_tokens`: both are SHARED from the target, and the head is consumed **packed** (2,542,796,800 B saved on device). Drafts on GB10 at k=7 with output token-identical to the target-only arm | The candidate selector is **not** quantized, contrary to that README — `candidate_selector.hidden_projection` is dense F16 and its two `[248320, 256]` codebooks are raw BF16, together ~17% of the file. Its `kernel_projection` and selector weights are **F16** where the surrounding norms are BF16, admitted by name only. The shipped **paged** draft route faults on run 2 in one process ([#2274](https://github.com/mudler/vllm.cpp/issues/2274)); `VT_DFLASH_PAGED=0` is the measured arm | +| Qwen3.8-27B DFlash2 draft, EXL3 5.0bpw (the drafter for the row above) | `model.safetensors` | 1,470,916,078 bytes; 189 tensors | `Mia-AiLab/Qwen3.8-27B-DFlash2-EXL3-5.0bpw` @ `4f0436269bca761b071f05319e8e04a87cc633f9` | `6b2e3afc694a343b7f3f0edfe5925e460762fc9ede4699165b577ca0733c8e56` | `DFlash2DraftModel`: five sliding-attention layers, `block_size 8`, target taps `[5,19,33,47,61]`. All 36 trellis modules use 5-bit codebook 2. The draft shares the target’s embeddings and packed output head. GB10 greedy output matches target-only generation at k=7 ([evidence](benchmarks/qwen38-27b-exl3-gb10.md)). | The candidate selector is unquantized: its projection is F16 and its two `[248320, 256]` codebooks are BF16. The loader admits F16 selector and convolution projections by name. Use `VT_DFLASH_PAGED=0` for the measured route; the paged route has a recorded repeat-run fault ([#2274](https://github.com/mudler/vllm.cpp/issues/2274)). | | Nemotron-3.5-Lightning-30B | `model-000{01..52}-of-00052.safetensors` | 21,583,809,748 bytes total | `nvidia/NVIDIA-Nemotron-3.5-Lightning-30B-A3B-NVFP4` @ `29f2d1746d8f41e316523194b19018707749b1b1` | `672c8bda10fdec0256e0819e112d2aa3a936cc3e5d311a05fd3ff773ca9a44b9` (first shard) | Device bf16, GQA, NVFP4 experts, and the NVFP4 head (A2-Q2b, unmeasured); host FP8 Mamba2 | GGUF, MTP, and batched decode | | MiniMax-H3 FL2VA | `MiniMax-H3-FL2VA-Q4_K_M.gguf` | 19,864,208,160 bytes | `realrebelai/MiniMax-H3_GGUFs` @ `daf03b4ca652cce16dfd4fcf91e79c52ffa5c1e7` | `5e8fa6e960d5fbd547390ceec63fcead275435d8f3bd2466a8a2cbd8c2e361e3` | Q4_K_M `t2va` and `fl2va`, verified end to end | `ref2va` requires the REF2VA partition | | MiniMax-H3 REF2VA | `MiniMax-H3-REF2VA-Q4_K_M.gguf` | 19,864,208,064 bytes | `realrebelai/MiniMax-H3_GGUFs` @ `daf03b4ca652cce16dfd4fcf91e79c52ffa5c1e7` | `17925612821ea3037ffaf5f7f9789f5460e87025385bd45e9ec6c7d536684d56` | Q4_K_M `ref2va`, verified end to end | `t2va` and `fl2va` require the FL2VA partition | @@ -1014,9 +1014,9 @@ repository in this project's history. | Qwen3.8-27B ModelOpt NVFP4 shard 3 of 4 | `model-00003-of-00004.safetensors` | 1,120,886,516 bytes | `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` @ `36f717a22990e82c54c1d48ee77c491b87825680` | Locally computed hash is owed; #821 | Same arms as shard 1 | The declared FP8 KV cache is unread; #1593 | | Qwen3.8-27B ModelOpt MTP drafter | `model-00004-of-00004.safetensors` | 849,400,592 bytes | `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` @ `36f717a22990e82c54c1d48ee77c491b87825680` | Locally computed hash is owed; #821 | Fifteen BF16 MTP tensors are present and unquantized | MTP execution is owed | | Qwen3.8-2.4T-A95B | `UD-Q1_0` ten-file GGUF split | about 370 GiB | `unsloth/Qwen3.8-2.4T-A95B-GGUF` @ `567d3e6ac26c5474b18311e619c04350fb9a5556` | `b7770552b2ac24e7334c917bc92e90e218e87cfe29484db65e62e8ef2a60334d` (shard 1); `2765517f833c736338d3ab34354e1c10eb8d79e62325f998285b435e5cf03dcd` (shard 2) | CPU expert streaming from disk | CUDA refuses a checkpoint that exceeds device capacity | -| Llama-3.2-1B-Instruct EXL3 3.0bpw (the first EXL3 checkpoint that GENERATES) | `model.safetensors` | 1,089,087,416 bytes | `turboderp/Llama-3.2-1B-Instruct-exl3` @ `f8f438c290680b15622270eff03bef23a458b1cf` (revision `3.0bpw` -- this repo publishes ONE BRANCH PER BIT WIDTH and `main` carries no weights at all, so a bare repo id resolves to nothing) | `3c0341e9c7c4c16a86a499de1dff4f6d7de9855541d669f3b0e214d72b54c2fc` | LOADS and GENERATES end to end through `vllm-cli` on `--device cpu`: `The capital of France is` -> ` Paris. Paris is known for its famous landmarks such as the Eiffel Tower` (greedy, 16 tokens, 2026-08-28). Native exllamav3 layout, no `.rank{r}` slicing; the body is 3-bit and `lm_head` is SIX-bit, resolved per tensor | Codebook **0** (the original QTIP 3INST), because the artifact ships no `mcg` marker and `LinearEXL3` derives the codebook from tensor PRESENCE. **It now RUNS ON CUDA**: the device arm instantiates `(3,0)`, `(3,1)` and `(6,0)`, so this checkpoint's 3-bit body and 6-bit head both reach the GPU. **It now RUNS ON ROCm too**, on `strix:gpu0` (`gfx1151`, RDNA3.5 APU, ROCm 7.2.4, 2026-09-02): the same file completes greedy generation with ZERO CPU reference-tier operations at 8.27 tok/s warm, against 0.83 tok/s for the same tree with the two ROCm registrations disabled, and the two arms emit the IDENTICAL continuation ` Paris. Paris is known for its famous` ([#2433](https://github.com/mudler/vllm.cpp/issues/2433)). That is an EXL3-vs-EXL3 A/B and NOT a ratio against a bf16 target: the BF16 control hung the GPU in the same lease ([#2511](https://github.com/mudler/vllm.cpp/issues/2511)), so this row records no BF16 denominator and no AMD clock attribution. gfx1151 is an APU with unified memory, so the figure generalizes to no discrete Radeon -- and on a discrete board the reference tier is off by design, which means this arm is what makes EXL3 RUN there at all rather than what makes it fast. It has NO GEMV fast path at `m == 1` and takes the regular shape table, which is upstream's behaviour too — its envelope refuses `bits != 4 && cb == 0` and its instantiation list omits `(3,0)`. **No speed is claimed.** The device figure ranged 2.1-5.0 tok/s across jobs on one binary and box, and an interleaved comparison against the BF16 twin of the same model read 2.33 vs 2.24 tok/s — indistinguishable, so the remaining ~50x to the memory floor is per-step engine overhead and not this scheme ([#2233](https://github.com/mudler/vllm.cpp/issues/2233)). q/k/v and gate/up run as separate GEMMs rather than one merged operand | -| DeepSeek-V4-Flash EXL3 trellis shard 1 of 172 | `exl3-layer-000-tp4-rank0.safetensors` | 515,850,920 bytes | `0xSero/deepseek-v4-flash-0731-spark` @ `22f28d32b9b29b4352eaa380ff8c2c170b2847ab` | `2ed7ae798a794019810b027fe2609e2cf4ad78d70b49c47b2970d03a0a7aaadf` | The rank-sliced EXL3 routed-expert tower LOADS (TP4 coalesced to TP1) and its experts EXECUTE through `vt::Exl3Gemm` on a CPU queue | The CUDA arm compiles for `sm_121a` and its numeric gates PASSED on GB10 on 2026-08-28 (`had_r_128` byte-identical, `exl3_gemm` `rel_rms 5.538e-4`, GEMV tier 3c `5.160e-4`); the FUSED MoE device arm still cannot run, because it needs a device-resident tower, so the routed experts execute on a CPU queue. That run decoded ZERO tensors of THIS artifact -- it found no readable shard -- so nothing here is a claim about these weights on a device. A SYNTHETIC rank-sliced checkpoint now loads and emits logits end to end; THIS artifact still does not, because its DSA compressor and indexer tensors are stored at twice the width the host forward indexes (`compressor.wgate` `[2*head_dim, H]`) and the loader refuses them BY NAME, and because its tokenizer is not read ([#1924](https://github.com/mudler/vllm.cpp/issues/1924)) | -| DeepSeek-V4-Flash EXL3 carried tower shard 1 of 5 | `carried-001.safetensors` | 4,288,630,252 bytes | `0xSero/deepseek-v4-flash-0731-spark` @ `22f28d32b9b29b4352eaa380ff8c2c170b2847ab` | `3b67ae29f1e75c2ecadfcafd3b0eecec640b06fd60b832f77e6bd3c2a8c85ccf` | The un-requantized `deepseek_v4_fp8` attention, router, shared-expert, compressor and embedding tensors, MATERIALIZED at load into the host-float tower the forward composes with — block-wise FP8 (`F8_E4M3` + `F8_E8M0` over 128x128 blocks) decoded to f32, BF16 norms and embeddings widened, I64 `tid2eid` narrowed to int32 | The DSA compressor and indexer tensors of this artifact are `2 * head_dim` / `2 * index_head_dim` wide and the loader refuses them by name (41 of its 43 layers carry a compressor); the 3,985 `mtp.*` NVFP4 draft tensors are skipped and counted, never silently dropped | +| Llama-3.2-1B-Instruct EXL3 3.0bpw (the first EXL3 checkpoint that GENERATES) | `model.safetensors` | 1,089,087,416 bytes | `turboderp/Llama-3.2-1B-Instruct-exl3` @ `f8f438c290680b15622270eff03bef23a458b1cf` (revision `3.0bpw` -- this repo publishes ONE BRANCH PER BIT WIDTH and `main` carries no weights at all, so a bare repo id resolves to nothing) | `3c0341e9c7c4c16a86a499de1dff4f6d7de9855541d669f3b0e214d72b54c2fc` | Loads through the shared dense path and generates with `vllm-cli --device cpu`, CUDA, and ROCm gfx1151. The native exllamav3 layout needs no rank slicing. Per-tensor metadata selects codebook 0, a 3-bit body, and a 6-bit output head. | CUDA GEMM supports both widths. The 3-bit codebook-0 body has no CUDA GEMV arm and uses regular GEMM. ROCm generation is verified on gfx1151. AMD clock attribution and discrete-GPU validation remain unavailable. See [ROCm evidence](../.agents/specs/backend-rocm-exl3.md) and [dense EXL3 limits](../.agents/specs/quant-exl3-shared.md). | +| DeepSeek-V4-Flash EXL3 trellis shard 1 of 172 | `exl3-layer-000-tp4-rank0.safetensors` | 515,850,920 bytes | `0xSero/deepseek-v4-flash-0731-spark` @ `22f28d32b9b29b4352eaa380ff8c2c170b2847ab` | `2ed7ae798a794019810b027fe2609e2cf4ad78d70b49c47b2970d03a0a7aaadf` | The loader coalesces the rank-sliced routed experts from TP4 to TP1. A synthetic checkpoint executes end to end; this does not establish support for the real artifact. | The real-artifact end-to-end gate remains open. CUDA GEMM/GEMV component gates do not verify the device-resident fused MoE tower. The loader now accepts the doubled compressor widths, and tokenizer parity is recorded. See the [current DeepSeek-V4 EXL3 blockers](../.agents/specs/model-dsv4-exl3.md#now). | +| DeepSeek-V4-Flash EXL3 carried tower shard 1 of 5 | `carried-001.safetensors` | 4,288,630,252 bytes | `0xSero/deepseek-v4-flash-0731-spark` @ `22f28d32b9b29b4352eaa380ff8c2c170b2847ab` | `3b67ae29f1e75c2ecadfcafd3b0eecec640b06fd60b832f77e6bd3c2a8c85ccf` | The loader materializes carried FP8, BF16, and F32 attention, router, shared-expert, compressor, and embedding weights into the host-float tower. It decodes block-wise FP8 (`F8_E4M3` with `F8_E8M0` scales over 128x128 blocks), widens BF16, and narrows I64 `tid2eid` to int32. | Compressor widths follow each layer’s compression ratio; the indexer compressor uses twice `index_head_dim`. Loading those tensors does not establish a real-artifact forward gate. The 3,985 `mtp.*` NVFP4 draft tensors remain skipped and counted. See [DeepSeek-V4 EXL3 status](../.agents/specs/model-dsv4-exl3.md#now). | | GLM-5.3-Flash FP8 source | `model-000{01..62}-of-00062.safetensors` | 328,326,771,576 bytes total (305.78 GiB) | `zai-org/GLM-5.3-Flash` @ `main`, read 2026-08-26 | Owed: no byte of payload has been fetched, so no local hash exists to state, and an unauthenticated tree hash is not a pin here | Declared source of `scripts/convert-glm5-next-gguf.py`. Only the safetensors HEADERS were read, by HTTP RANGE over all 62 shards: 76,108 tensors, `F8_E4M3` block-quantized at `weight_block_size: [128, 128]` with `weight_scale_inv` companions, plus BF16 and F32 scales | **Nothing has been converted.** The download needs explicit developer authority and a box with room for 305.78 GiB of source and ~100.35 GiB of output at once; owed as O7 on [#2011](https://github.com/mudler/vllm.cpp/issues/2011). The revision is a branch name and not a commit, which is NOT a pin: it is what was read, and W7b re-reads and records the commit when it stages the bytes | | GLM-5.3-Flash GGUF | `GLM-5.3-Flash-UD-Q2_K_XL-0000{1..4}-of-00004.gguf` | 108,720,071,427 bytes total (101.2535 GiB) across four shards; 1412 tensors | `unsloth/GLM-5.3-Flash-GGUF` @ `d425e572fb9686125831f476129e51cea34bc5b4`, path `UD-Q2_K_XL`, staged 2026-08-28 | Owed for this row: the shards are staged and were sha256-verified when they were fetched, but **W5c consumed only the four GGUF HEADERS** and states no hash of its own. W7b ([#2225](https://github.com/mudler/vllm.cpp/issues/2225)) records the per-shard sha256 alongside the load it measures | **LOADS on `--device cpu`, and the engine's multi-KV guard no longer refuses above the model's forward** ([#2348](https://github.com/mudler/vllm.cpp/issues/2348)). **A MATERIALIZED LOAD EXISTS** -- driven at this artifact on `dgx:gpu0` 2026-08-30, all four shards load and the engine sizes its caches in under 26 minutes wall ([#2343](https://github.com/mudler/vllm.cpp/issues/2343)). At that change the first step threw at the `multi_kv` guard above the model's own hook; W5b-2c ([#2348](https://github.com/mudler/vllm.cpp/issues/2348)) writes the consuming forward that guard was waiting for and it no longer fires for this model. **THIS ARTIFACT GENERATES COHERENT TEXT, and peak RSS is MEASURED** as of [#2241](https://github.com/mudler/vllm.cpp/issues/2241). On `dgx:gpu0` 2026-08-30, in the SHIPPED configuration with no diagnostic env set, `vllm-cli --device cpu --max-tokens 2` at the prompt `The capital of France is` emits ` Paris.` at `rc=0`, and `VmHWM` peaks at 104,792,300 kB = 99.94 GiB. Two instrumented `thor:gpu0` runs the same day supply the bisect: four tokens read ` Paris. Paris is`, the prefill top-5 is ` Paris` (16.427), ` one`, ` located`, ` known`, ` a` at margin 1.279, and none of 180 per-layer readings over four steps carries a NaN. The first attempt emitted token id 0 eight times, because the loader repacked this file's 346 q8_0 tensors into the i8mm interleave that the host bridge reads as plain blocks (spec `## Owed` O30). **No speed number is claimed, and the earlier ones are void** -- they were taken from an all-NaN forward. The GB10 arm is the one measured above. The GGUF arm of `load_weights` resolves all 1383 backbone tensors of this file (W5c, [#2242](https://github.com/mudler/vllm.cpp/issues/2242)); `blk.45`, the multi-token-prediction block, is read, counted and DROPPED, as the transformers reference does. `ModelRegistry::Forward` dispatches to the model as of W5b-2b ([#2337](https://github.com/mudler/vllm.cpp/issues/2337)), which bridges ONE decoder layer at a time out of the block-resident tower and decodes only the 8 of 288 experts a token selects — a float tower is 426.72 GiB against ~119.63 GiB usable. **A MATERIALIZED LOAD NOW EXISTS**: driven at this artifact on `dgx:gpu0` 2026-08-30, all four shards load and the engine sizes its caches in under 26 minutes wall. **NO TOKEN WAS GENERATED** — the first step throws at the `multi_kv` guard above the model's own hook ([#2343](https://github.com/mudler/vllm.cpp/issues/2343), [#2068](https://github.com/mudler/vllm.cpp/issues/2068)) — and **peak RSS and speed are still unmeasured**, because the staging run did not sample them. The vision tower (a separate `mmproj-BF16.gguf`) and the safetensors arm still refuse by name, as does a multi-request step; a non-CPU queue is admitted as of W9c-3a ([#2464](https://github.com/mudler/vllm.cpp/issues/2464)) for the routed-expert GEMM alone, and a device that is neither CPU nor CUDA is refused by name; **the KV-cache spec does not**, as of W5 ([#2223](https://github.com/mudler/vllm.cpp/issues/2223)), which publishes its three groups through the production factory hook | **The earlier row here said `none exists`, and that was true when it was written (2026-08-26) and is not now.** "UD-Q2_K_XL" names a TARGET AVERAGE and not a format: the census over all 1412 tensors is F32 638, Q8_0 346, Q5_K 181, Q6_K 117, IQ2_XS 82, IQ3_XXS 41, IQ4_XS 3, Q2_K 2, Q4_K 1, Q3_K 1 — **two** Q2_K tensors in a file named Q2_K. It fits `dgx:gpu0` only because IQ2_XS and IQ4_XS keep their blocks ([#2247](https://github.com/mudler/vllm.cpp/issues/2247)); both now have a CUDA keep-quant kernel too ([#2260](https://github.com/mudler/vllm.cpp/issues/2260)), so the expert GEMM no longer drains the stream to the host and the fused seam no longer throws. W9c-3a ([#2464](https://github.com/mudler/vllm.cpp/issues/2464)) then built a device arm for this artifact's routed-expert GEMM and MEASURED it end to end, where it **SEGFAULTED**: both `--device cuda` legs on `dgx:gpu0` died with rc=139 emitting no token, reproducibly (spec O46). The split is therefore OPT-IN and defaults OFF, so `--device cuda` refuses exactly as it did before. **Use `--device cpu`** -- measured on that artifact it emits ` Paris.` at rc=0, 1176 s wall of which 169 s is generation. Every OTHER primitive of this model is still a host reference on an interposed CPU queue (spec O43), so what `--device cuda` reaches is one arm of eleven and not a device arm. **A materialized load NOW exists and a token still does not** — `dgx:gpu0` 2026-08-30 ([#2343](https://github.com/mudler/vllm.cpp/issues/2343)): all four shards load and the engine sizes its caches, then the first step throws at the `multi_kv` guard above the model's own hook. **Peak RSS and speed remain unmeasured** | | GLM-5.3-Flash config | `config.json` | 69,416 bytes | `zai-org/GLM-5.3-Flash` @ `main`, read 2026-08-27 | sha256 `bb8f01c42cb92a52ca72e65afb4d5bd8d11aef083cd210e8de25dfb904f23e9f` | The ONLY byte of this checkpoint any change on this row has consumed. Checked in verbatim as `tests/vllm/models/fixtures/glm5_next/config.json` and used as W1's gate fixture, so the config layer is gated against what the checkpoint says rather than against what a port's author believed it says | **Arms refused by name:** the SAFETENSORS one, which is what this row is, because every published safetensors artifact of this model exceeds every device this project owns. `Glm5NextForConditionalGeneration` is REGISTERED, its config RESOLVES, and the GGUF arm both loads and forwards ([#2067](https://github.com/mudler/vllm.cpp/issues/2067), [#2242](https://github.com/mudler/vllm.cpp/issues/2242), [#2337](https://github.com/mudler/vllm.cpp/issues/2337)). The revision is a branch name and not a commit, which is NOT a pin for the WEIGHTS; for this one file the sha256 above is the pin |