Unsloth Dynamic 3.0 model files

The engine now runs unsloth's Qwen3.8-Flash-Next quants (UD-IQ1_S up to UD-Q6_K_XL and Q8_0) as native packs, besides the ISTA-DASLab files. - Split models: --native takes the first shard and finds the others by name (gguf_split_paths, GgufModel). The expert arena is filled from the shards by tensor name and checked against native_experts.txt, whose lines no longer carry shard-1 offsets. The native head and embedding are found in any shard, and --ple-gguf defaults to the shard that holds per_layer_token_embd. - GPU kernels: expert dot products and dequantizers for Q4_K, Q5_K, Q6_K, Q5_1, Q8_0, IQ1_S and IQ4_XS (llama.cpp's), one dispatch list for all formats, and a startup check that every layer's formats have kernels. The Q5_1 min and IQ1_S delta terms use the quantized activation sums, as ggml-cpu does: 1.1% instead of 1.9% relative error per expert. - PLE: the table may be IQ4_NL or Q8_0 and may share its shard with other tensors; the native PLE key may be any native MMVQ type (Q8_0 here). - tools/iq_pack.py: all shards of a split model, and every float tensor in the form the engine takes. The files' F32 router, SSM gate and injection tensors hold BF16 values (exact); their Q8_0 hyper-connection and PLE value weights are re-rounded to BF16. - native_expert_parity: split models, a --synthetic mode for formats with no model file at hand, and a check of the GPU dequantizers. - README: the unsloth files, their RAM needs and how to pack them. UD-Q4_K_XL on an RTX 3090 (160 GB RAM): prompts 529-594 tok/s, generation 49-57 tok/s. Greedy tokens equal llama.cpp's argmax on the same GGUF at 98.2% (code) and 95.6% (thinking) of positions.
2026-11-07 22:59:00 -10:00 · 2026-09-25 01:54:29 +02:00
1 parent 0aaa754656
commit 8029fa9df7
@@ -74,6 +74,28 @@ of [Qwen3.8-Flash-Next](https://huggingface.co/Qwen/Qwen3.8-Flash-Next).

 With 64 GB of RAM all three fit (close the browser for IQ3_XXS, and keep its context at 128K or less). With 48 GB only Q2_0 / IQ2_XS may fit. 32 GB is not enough.

+### unsloth's Unsloth Dynamic 3.0 files
+
+The engine also runs [unsloth's quantizations](https://huggingface.co/unsloth/Qwen3.8-Flash-Next-GGUF), which
+START-HERE does not set up. All experts sit in RAM, pinned when they fit the driver's limit:
+
+| Model | Download | Experts in RAM | n-gram table |
+| --- | ---: | ---: | --- |
+| UD-IQ1_S / UD-IQ1_M | 73-75 GB | 40-42 GB | IQ4_NL, 28.8 GB |
+| UD-Q2_K_XL / UD-IQ3_XXS | 79-82 GB | 46-49 GB | IQ4_NL |
+| UD-Q3_K_XL / UD-IQ4_XS | 90-94 GB | 56-60 GB | IQ4_NL |
+| UD-Q4_K_XL | 111 GB | 77 GB | IQ4_NL |
+| UD-Q5_K_XL / UD-Q6_K_XL | 158-169 GB | 98-109 GB | Q8_0, 54.4 GB |
+| Q8_0 | 188 GB | 128 GB | Q8_0 |
+
+Pack a model once (seconds: the tokenizer and the float tensors; the experts are read from the GGUF at start):
+
+    python tools/iq_pack.py --gguf <dir>/<model>-00001-of-0000N.gguf --out packs/<name>
+
+and start it with `--pack packs/<name> --native <dir>/<model>-00001-of-0000N.gguf` plus the usual flags (`--spec`,
+`--mtp`, `--prefill`, ...). The other shards, and the one holding the n-gram table, are found by name.
+UD-Q4_K_XL on an RTX 3090 with 160 GB of RAM: prompts 529-594 tok/s, generation 49-57 tok/s.
+
 ## Before you start

 You need **only an NVIDIA driver** (version 580 or newer; update it with the NVIDIA App or from
@@ -23,12 +23,15 @@
 #include <cstdint>
 #include <cstdio>
 #include <cstring>
+#include <memory>
 #include <string>
 #include <vector>
 #include <map>
 #include <stdexcept>
 #include <algorithm>

+#include "strata/artifact/gguf_split.hpp"
+
 #ifdef _WIN32
 #define WIN32_LEAN_AND_MEAN
 #include <windows.h>
@@ -338,6 +341,7 @@ public:
    uint32_t version() const { return version_; }
    uint64_t data_start() const { return data_start_; }
    uint64_t file_size() const { return size_; }
+    uint64_t alignment() const { return alignment_; }

    const TensorInfo* find(const std::string& name) const {
        for (const auto& t : tensors_)
@@ -355,6 +359,7 @@ public:
        return it == meta_.end() ? nullptr : &it->second;
    }
    const uint8_t* tensor_data(const TensorInfo& t) const { return base_ + data_start_ + t.offset; }
+    const std::string& path() const { return path_; }

 private:
    void open() {
@@ -451,6 +456,31 @@ private:
    std::map<std::string, MetaValue> meta_;
 };

+// The opened shards of one model, tensors looked up across all of them.  Shard 0 holds the metadata.
+class GgufModel {
+public:
+    explicit GgufModel(const std::vector<std::string>& paths) {
+        if (paths.empty()) throw std::runtime_error("GGUF: a model needs at least one shard");
+        for (const auto& p : paths) shards_.push_back(std::make_unique<GgufFile>(p));
+    }
+    size_t size() const { return shards_.size(); }
+    const GgufFile& shard(size_t i) const { return *shards_[i]; }
+    /// The tensor named `name` (its shard index in `*shard`), or nullptr.  Throws when two shards hold it.
+    const TensorInfo* find(const std::string& name, size_t* shard = nullptr) const {
+        const TensorInfo* found = nullptr;
+        for (size_t i = 0; i < shards_.size(); ++i)
+            if (const TensorInfo* t = shards_[i]->find(name)) {
+                if (found) throw std::runtime_error("GGUF: tensor " + name + " is in two shards");
+                found = t;
+                if (shard) *shard = i;
+            }
+        return found;
+    }
+
+private:
+    std::vector<std::unique_ptr<GgufFile>> shards_;
+};
+
 // ---- architecture guard (P1.S2). The engine is specialised to ONE model; anything else must be
 // refused with a precise error rather than silently mis-run.
 struct Qwen4ExpGuard {
@@ -0,0 +1,38 @@
+// include/strata/artifact/gguf_split.hpp - the shards of a split model.  Separate from gguf_reader.hpp, which
+// includes <windows.h>, so any translation unit can resolve a model's shards.
+#pragma once
+
+#include <cstdio>
+#include <fstream>
+#include <stdexcept>
+#include <string>
+#include <vector>
+
+namespace strata {
+
+// Split models: `<name>-00001-of-0000N.gguf` ... `<name>-0000N-of-0000N.gguf`, as llama.cpp's gguf-split
+// names them.  The shards in split order, from the first one's path; a file without that suffix is a model of
+// one shard.  Throws when a shard is missing.
+inline std::vector<std::string> gguf_split_paths(const std::string& first) {
+    static const char pat[] = "-00001-of-";
+    const size_t at = first.rfind(pat);
+    const size_t digits = at == std::string::npos ? 0 : at + sizeof(pat) - 1;
+    if (at == std::string::npos || first.size() != digits + 10 || first.compare(digits + 5, 5, ".gguf") != 0)
+        return {first};
+    int n = 0;
+    for (size_t i = digits; i < digits + 5; ++i) {
+        if (first[i] < '0' || first[i] > '9') return {first};
+        n = n * 10 + (first[i] - '0');
+    }
+    std::vector<std::string> out;
+    for (int i = 1; i <= n; ++i) {
+        char no[8];
+        std::snprintf(no, sizeof no, "%05d", i);
+        std::string p = first.substr(0, at + 1) + no + first.substr(at + 6);
+        if (!std::ifstream(p, std::ios::binary)) throw std::runtime_error("missing model shard " + p);
+        out.push_back(std::move(p));
+    }
+    return out;
+}
+
+}  // namespace strata
@@ -321,8 +321,9 @@ public:
    /// Allocates and loads `<pack_dir>/experts.bin`.  Prints nothing; the caller reports `note()` and the load
    /// rate, because those are the two numbers that say whether the arena is the one that was asked for.
    bool open(const std::string& pack_dir, int64_t n_layers, int64_t n_expert, int threads, std::string& err);
-    /// Plan v0.3 P6: a native pack without experts.bin takes its experts from the model's shard 1.
-    void set_gguf(const std::string& shard1) { gguf_ = shard1; }
+    /// Plan v0.3 P6: a native pack without experts.bin takes its experts from the model's GGUF shards
+    /// (strata::gguf_split_paths), found by tensor name.
+    void set_gguf(const std::vector<std::string>& shards) { gguf_ = shards; }
    void close();

    bool mapped() const { return base_ != nullptr; }
@@ -348,7 +349,7 @@ private:
    std::string note_;
    double gib_per_s_ = 0.0;
    uint64_t pinned_bytes_ = 0;
-    std::string gguf_;
+    std::vector<std::string> gguf_;
 };

 }  // namespace strata::core
@@ -3,12 +3,14 @@
 #include <cstddef>
 #include <cstdint>
 #include <string>
+#include <vector>

 namespace strata::core {

 /// Opt-in native output projection. Uploads unchanged GGUF blocks once at startup and owns its
 /// activation scratch. The normal canonical head remains available for numerical A/B comparisons.
-/// Plan v0.3 P6: any type the native MMVQ takes (Q5_K in the Q2_0 / IQ3_XXS files, IQ4_XS in IQ2_XS).
+/// Plan v0.3 P6: any type the native MMVQ takes (Q5_K in the Q2_0 / IQ3_XXS files, IQ4_XS in IQ2_XS, Q8_0 in
+/// UD-Q4_K_XL).  `load` takes the model's shards (strata::gguf_split_paths).
 class NativeHead {
 public:
    NativeHead() = default;
@@ -16,7 +18,7 @@ public:
    NativeHead(const NativeHead&) = delete;
    NativeHead& operator=(const NativeHead&) = delete;

-    bool load(const std::string& gguf, int64_t n_in, int64_t n_out, std::string& err);
+    bool load(const std::vector<std::string>& shards, int64_t n_in, int64_t n_out, std::string& err);
    bool run(const float* mixed, float* logits, void* stream, std::string& err) const;
    uint64_t weight_bytes() const { return bytes_; }
    bool loaded() const { return weights_ != nullptr; }
@@ -43,7 +45,7 @@ public:
    ~NativeEmbed();
    NativeEmbed(const NativeEmbed&) = delete;
    NativeEmbed& operator=(const NativeEmbed&) = delete;
-    bool load(const std::string& gguf, int64_t n_embd, int64_t n_vocab, std::string& err);
+    bool load(const std::vector<std::string>& shards, int64_t n_embd, int64_t n_vocab, std::string& err);
    /// Rows for device token ids.
    void gather_dev(const int32_t* tokens, int64_t n_tok, float* out, void* stream) const;
    /// One row for a host token id.
@@ -20,9 +20,6 @@ struct ExpertLayout {
    int64_t n_layers = 0, n_expert = NE;
    std::vector<NativeFmt> fmt;           ///< per layer (native packs)
    std::vector<uint64_t> offset, bytes;  ///< per layer: where its 512 blobs start, bytes per blob
-    /// Plan v0.3 P6: per layer, the absolute offsets of the gate / up / down tensors in the model's shard 1, so
-    /// the arena can be filled from the GGUF itself when the pack has no experts.bin (3 x n_layers, 0 = unknown).
-    std::vector<uint64_t> gguf_off;
    uint64_t max_blob = BLOB;
    uint64_t total = 0;                   ///< experts.bin size

@@ -1,5 +1,6 @@
-// include/strata/kernels/iq_kernels.hpp - the i-quant formats (IQ1_M, IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ3_S,
-// IQ4_NL) and Q2_0 on the GPU for the IQ2_XS / IQ3_XXS model files.
+// include/strata/kernels/iq_kernels.hpp - the GGUF expert and embedding formats on the GPU: the i-quants (IQ1_M,
+// IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ3_S, IQ4_NL) and Q2_0 of the IQ2_XS / IQ3_XXS model files, and Q4_K, Q5_K,
+// Q5_1 and Q8_0 of the UD-Q4_K_XL file.
 //
 // The block layouts, codebook grids and dot products are llama.cpp's (ggml-common.h, ggml-cuda/vecdotq.cuh,
 // ggml-cuda/dequantize.cuh; MIT, see third_party/ggml/LICENSE and VERSION.txt), so a weight means exactly what it
@@ -43,6 +44,9 @@ struct NativeExpertLayout {
    size_t bytes = 0;                   // the whole blob
 };
 NativeExpertLayout native_expert_layout(int gu_type, int d_type, int64_t n_embd, int64_t n_ff);
+/// Whether the GPU has the dot products (native_expert_grouped) and the dequantizers (the prompt path) for a layer's
+/// gate/up and down formats at this geometry.
+bool native_expert_supported(int gu_type, int d_type, int64_t n_embd, int64_t n_ff) noexcept;

 /// Bytes of scratch `native_expert_grouped` needs for `cap_entries` entries.
 size_t native_expert_scratch_bytes(int64_t cap_entries, int64_t n_ff);
@@ -5,9 +5,9 @@
 //   * `ngram_rows` - the HOST hash.  Sixteen row indices per token from the last three token ids, computed
 //     with 64-bit multiply/xor and no tensor op at all.  `qwen4exp.cpp` L1092-1124, where the source says
 //     outright that it is host-side because "ggml has no int64 and no xor".
-//   * `PleTable`  - the row gather.  `per_layer_token_embd.weight` is IQ4_NL and is NOT IN THE PACK: it is
-//     51.2e9 elements (28.8 GB) in the ORIGINAL second GGUF shard, and it is the only tensor this engine
-//     reads from the GGUF rather than from the canonical pack.
+//   * `PleTable`  - the row gather.  `per_layer_token_embd.weight` is IQ4_NL (Q8_0 in the unsloth UD-Q5_K_XL and
+//     larger files) and is NOT IN THE PACK: it is 51.2e9 elements (28.8 GB; 54.4 GB as Q8_0) in the model's
+//     GGUF, and the engine reads it from there.
 //
 // THE HASH IS THE ONE PLACE IN THIS MODEL WHERE A WRONG BIT IS SILENTLY PLAUSIBLE.  Every one of its
 // properties has a rival reading that produces a valid index in range:
@@ -24,6 +24,7 @@

 #include <cstdint>
 #include <string>
+#include <vector>

 namespace strata::kernels {

@@ -45,6 +46,7 @@ inline constexpr float NG_RMS_EPS = 1e-6f;
 // 5 blocks of 32 at 18 bytes = 90 bytes.  The head-slowest flatten then makes 16 rows exactly n_embd = 2560.
 inline constexpr uint64_t PLE_TABLE_ROWS = 320001536ull;
 inline constexpr int PLE_ROW_BYTES = (PLE_HEAD_DIM / 32) * 18;           // 90
+inline constexpr int PLE_ROW_BYTES_Q8_0 = (PLE_HEAD_DIM / 32) * 34;      // 170, the Q8_0 tables

 /// The artifact's own hash constants, transcribed from `docs/gguf-dump-shard1.txt`:
 ///
@@ -95,6 +97,8 @@ int iq4nl_code(int code);
 /// so byte j carries elements j and j+16 - NOT 2j and 2j+1, which is the natural reading and which produces
 /// a perfectly plausible embedding of the wrong 160 values.
 void iq4nl_dequant_row(const uint8_t* row, float* out160);
+/// One row of a table of GGML type `type` (IQ4_NL = 20 or Q8_0 = 8) -> 160 floats.
+void ple_dequant_row(int type, const uint8_t* row, float* out160);

 /// How the table's rows are read (plan v0.3 P2). `Direct` is the default: unbuffered 4 KiB reads from the SSD,
 /// so the table never occupies RAM or the OS file cache. `Mmap` is the earlier memory-mapped path, kept as the
@@ -112,6 +116,9 @@ struct PleIoOptions {
    bool io_thread = true;           ///< reads submitted by a worker thread, not the caller
 };

+/// The shard among `shards` that holds `per_layer_token_embd.weight`, or "" (headers only).
+std::string ple_table_shard(const std::vector<std::string>& shards);
+
 /// The PLE table.  Held by pointer-to-impl so this header does not drag `<windows.h>` into every
 /// translation unit that wants the hash.
 class PleTable {
@@ -1,6 +1,7 @@
 // include/strata/ngram/ple_reader.hpp - plan v0.3 P2: the n-gram table read straight from the SSD.
 //
-// The table is 320,001,536 rows of 90 bytes (26.8 GiB) and is never held in RAM: every row comes from an
+// The table is 320,001,536 rows of 90 bytes (IQ4_NL, 26.8 GiB; 170 bytes as Q8_0) and is never held in RAM: every
+// row comes from an
 // unbuffered 4 KiB read (platform::DirectFile). A token needs 16 rows on 16 different pages, and all 16 depend
 // on the token itself, so the only time to hide them is the embedding plus layer 0. Hence the split API:
 //
@@ -11,7 +12,7 @@
 // `issue` also serves prefill: pass all 16 x N rows of a chunk; pages are deduplicated, sorted by offset and
 // kept at most `max_inflight` deep, so a chunk's reads can run while the previous chunk computes.
 //
-// THE ROW CACHE IS NOT THE TABLE. It keeps rows this process has already fetched (90 bytes each, bounded,
+// THE ROW CACHE IS NOT THE TABLE. It keeps rows this process has already fetched (one row each, bounded,
 // clock eviction). Measured on the frozen corpus (bench/results/2026-09-23-ngram-io): 1M rows (~95 MB)
 // would serve up to ~82% of reads; within one long prompt 20-34% of rows recur. Capacity 0 disables it.
 #pragma once
@@ -24,7 +25,7 @@

 namespace strata::ngram {

-inline constexpr uint32_t ROW_BYTES = 90;
+inline constexpr uint32_t ROW_BYTES = 90;       ///< an IQ4_NL row, the default
 inline constexpr uint32_t PAGE = 4096;

 struct ReaderStats {
@@ -52,17 +53,17 @@ public:
    PleReader(const PleReader&) = delete;
    PleReader& operator=(const PleReader&) = delete;

-    /// `table_offset` is the byte offset of row 0 in the file and `n_rows` the row count; both come from a
-    /// validated GGUF parse (PleTable::open checks the table exactly fills the file from there).
+    /// `table_offset` is the byte offset of row 0 in the file, `n_rows` the row count and `row_bytes` a row's
+    /// size; all come from a validated GGUF parse (PleTable::open checks the table exactly fills its span).
    /// `io_thread` (default): a worker thread submits and reaps reads, so `issue` costs the caller no ReadFile
    /// calls. false: the caller's thread does it (A/B arm).
    bool open(const std::string& path, uint64_t table_offset, uint64_t n_rows, uint32_t max_inflight,
-              uint64_t cache_rows, std::string& err, bool io_thread = true);
+              uint64_t cache_rows, std::string& err, bool io_thread = true, uint32_t row_bytes = ROW_BYTES);
    void close();
    bool is_open() const;

-    /// Start fetching `n` rows; row i's 90 raw bytes land at `out_raw + 90 * i`. `out_raw` must stay valid
-    /// until `collect` returns. Out-of-range rows produce 90 zero bytes (the mmap path's behaviour).
+    /// Start fetching `n` rows; row i's raw bytes land at `out_raw + row_bytes * i`. `out_raw` must stay valid
+    /// until `collect` returns. Out-of-range rows produce zero bytes (the mmap path's behaviour).
    Ticket issue(const uint32_t* rows, size_t n, uint8_t* out_raw);

    /// Block until every row of the ticket is in `out_raw`. Returns false on an I/O error (message in `err`).
@@ -1,5 +1,6 @@
 // src/core/expert_source.cpp - the adapter.  See the header for the three clauses of the contract.
 #include "strata/core/expert_source.hpp"
+#include "strata/artifact/gguf_reader.hpp"
 #include "strata/kernels/cpu/expert_layout.hpp"

 #include "strata/core/pinned.hpp"
@@ -535,19 +536,57 @@ void expert_hit_run(void* user, void* stream, HitPhase phase, const int32_t* ids

 // ================================ THE RESIDENT ARENA (R2.1) ================================

-// Plan v0.3 P6: the arena from the model's shard 1.  Each layer's gate, up and down tensors hold the 512 experts
-// one after another; they are read in chunks and each expert's slice lands at its place in the blob
+namespace {
+// Plan v0.3 P6: the arena straight from the model's GGUF shards.  Each layer's gate, up and down tensors hold the
+// 512 experts one after another; they are read in chunks and each expert's slice lands at its place in the blob
 // [gate rows | up rows | down rows] - the layout tools/iq_pack.py would have written to experts.bin.
-LoadStats load_experts_gguf(const std::string& gguf, uint8_t* dst, const strata::kernels::cpu::ExpertLayout& lay,
-                            int threads) {
+struct GgufSpan {
+    size_t shard = 0;
+    uint64_t offset = 0;   ///< in the shard's file
+};
+
+/// Where each layer's gate, up and down tensors are (3 per layer), found by name in whichever shard holds them
+/// and checked against the pack's formats, so a pack and a model that do not belong together are refused.
+bool locate_experts_gguf(const std::vector<std::string>& shards, const strata::kernels::cpu::ExpertLayout& lay,
+                         std::vector<GgufSpan>& out, std::string& err) {
+    static const char* roles[3] = {"gate", "up", "down"};
+    out.assign((size_t) (3 * lay.n_layers), GgufSpan{});
+    try {
+        const strata::GgufModel model(shards);
+        for (int64_t l = 0; l < lay.n_layers; ++l) {
+            const auto& fm = lay.fmt[(size_t) l];
+            const uint64_t per[3] = {fm.up_off, fm.up_off, lay.bytes[(size_t) l] - fm.down_off};
+            for (int r = 0; r < 3; ++r) {
+                const std::string name = "blk." + std::to_string(l) + ".ffn_" + roles[r] + "_exps.weight";
+                size_t s = 0;
+                const strata::TensorInfo* t = model.find(name, &s);
+                int be = 0, bb = 0;
+                if (t == nullptr || (int) t->type != (r < 2 ? fm.gu_type : fm.d_type) ||
+                    !strata::block_geometry(t->type, be, bb) ||
+                    t->elements() / (uint64_t) be * (uint64_t) bb != per[r] * (uint64_t) lay.n_expert) {
+                    err = "ArenaExpertSource: " + name + " is missing from the model or differs from the pack's "
+                          "native_experts.txt";
+                    return false;
+                }
+                out[(size_t) (3 * l + r)] = GgufSpan{s, model.shard(s).data_start() + t->offset};
+            }
+        }
+        return true;
+    } catch (const std::exception& e) {
+        err = std::string("ArenaExpertSource: ") + e.what();
+        return false;
+    }
+}
+
+LoadStats load_experts_gguf(const std::vector<std::string>& shards, const std::vector<GgufSpan>& spans, uint8_t* dst,
+                            const strata::kernels::cpu::ExpertLayout& lay, int threads) {
    LoadStats st;
    st.layers = (uint64_t) lay.n_layers;
    const auto t0 = std::chrono::steady_clock::now();
    std::atomic<int64_t> next{0};
    std::atomic<bool> bad{false};
    auto worker = [&]() {
-        std::ifstream f(gguf, std::ios::binary);
-        if (!f) { bad = true; return; }
+        std::vector<std::ifstream> files(shards.size());
        std::vector<uint8_t> buf;
        for (;;) {
            const int64_t l = next.fetch_add(1);
@@ -557,7 +596,11 @@ LoadStats load_experts_gguf(const std::string& gguf, uint8_t* dst, const strata:
            const uint64_t per[3] = {fm.up_off, fm.up_off, blob - fm.down_off};
            const uint64_t at[3] = {0, fm.up_off, fm.down_off};
            for (int r = 0; r < 3; ++r) {
-                const uint64_t src = lay.gguf_off[(size_t) (3 * l + r)];
+                const GgufSpan& sp = spans[(size_t) (3 * l + r)];
+                std::ifstream& f = files[sp.shard];
+                if (!f.is_open()) f.open(shards[sp.shard], std::ios::binary);
+                if (!f) { bad = true; return; }
+                const uint64_t src = sp.offset;
                const uint64_t total = per[r] * (uint64_t) lay.n_expert;
                const uint64_t chunk = per[r] * 16;           // 16 experts per read
                buf.resize((size_t) chunk);
@@ -586,6 +629,7 @@ LoadStats load_experts_gguf(const std::string& gguf, uint8_t* dst, const strata:
    st.seconds = std::chrono::duration<double>(std::chrono::steady_clock::now() - t0).count();
    return st;
 }
+}  // namespace

 ArenaExpertSource::~ArenaExpertSource() { close(); }

@@ -603,10 +647,13 @@ bool ArenaExpertSource::open(const std::string& pack_dir, int64_t n_layers, int6
    const uint64_t want = lay.total;

    // plan v0.3 P6: no experts.bin in a native pack -> the experts come straight from the GGUF
-    const bool from_gguf = !std::ifstream(path, std::ios::binary) && lay.native && !lay.gguf_off.empty() && !gguf_.empty();
+    const bool from_gguf = !std::ifstream(path, std::ios::binary) && lay.native && !gguf_.empty();
    // SIZE CHECK BEFORE THE ALLOCATION, not after.  A wrong pack should name the two numbers rather than spend
    // 34 GB and a minute of loading first.
-    if (!from_gguf) {
+    std::vector<GgufSpan> spans;
+    if (from_gguf) {
+        if (!locate_experts_gguf(gguf_, lay, spans, err)) return false;
+    } else {
        std::ifstream f(path, std::ios::binary | std::ios::ate);
        if (!f) { err = "ArenaExpertSource: cannot open " + path; return false; }
        const uint64_t got = (uint64_t) f.tellg();
@@ -637,7 +684,7 @@ bool ArenaExpertSource::open(const std::string& pack_dir, int64_t n_layers, int6
        err = "ArenaExpertSource: the arena could not be reserved (" + std::to_string(want) + " B)";
        return false;
    }
-    const LoadStats st = from_gguf ? load_experts_gguf(gguf_, a->data(), lay, threads)
+    const LoadStats st = from_gguf ? load_experts_gguf(gguf_, spans, a->data(), lay, threads)
                                   : load_experts_ranges(path, a->data(), loff, lbytes, threads, /*chunk=*/8u << 20);
    if (st.bytes != want) {
        delete a;
@@ -15,22 +15,19 @@ NativeHead::~NativeHead() {
    if (weights_) cudaFree(weights_);
 }

-bool NativeHead::load(const std::string& path, int64_t n_in, int64_t n_out, std::string& err) {
+bool NativeHead::load(const std::vector<std::string>& shards, int64_t n_in, int64_t n_out, std::string& err) {
    if (loaded()) { err = "native head is already loaded"; return false; }
    if (n_in <= 0 || n_out <= 0 || n_in > INT_MAX || n_out > INT_MAX || n_in % 256) {
        err = "native head requires positive int32 dimensions and whole 256-value rows";
        return false;
    }
    try {
-        strata::GgufFile gguf(path);
-        err = strata::check_architecture(gguf);
+        const strata::GgufModel model(shards);
+        err = strata::check_architecture(model.shard(0));
        if (!err.empty()) return false;
-        const strata::TensorInfo* tensor = nullptr;
-        for (const auto& candidate : gguf.tensors()) {
-            if (candidate.name != "output.weight") continue;
-            if (tensor) { err = "native head: duplicate output.weight"; return false; }
-            tensor = &candidate;
-        }
+        size_t at = 0;
+        const strata::TensorInfo* tensor = model.find("output.weight", &at);
+        const strata::GgufFile& gguf = model.shard(at);
        if (!tensor || !strata::kernels::native_mmvq_supported((int) tensor->type) || tensor->shape.size() != 2 ||
            tensor->shape[0] != (uint64_t) n_in || tensor->shape[1] != (uint64_t) n_out) {
            err = "native head: expected a natively supported output.weight with the canonical head dimensions";
@@ -104,12 +101,12 @@ NativeEmbed::~NativeEmbed() {
    if (host_) cudaFreeHost(host_);
 }

-bool NativeEmbed::load(const std::string& path, int64_t n_embd, int64_t n_vocab, std::string& err) {
+bool NativeEmbed::load(const std::vector<std::string>& shards, int64_t n_embd, int64_t n_vocab, std::string& err) {
    try {
-        strata::GgufFile gguf(path);
-        const strata::TensorInfo* t = nullptr;
-        for (const auto& c : gguf.tensors())
-            if (c.name == "token_embd.weight") t = &c;
+        const strata::GgufModel model(shards);
+        size_t at = 0;
+        const strata::TensorInfo* t = model.find("token_embd.weight", &at);
+        const strata::GgufFile& gguf = model.shard(at);
        if (!t || t->shape.size() != 2 || t->shape[0] != (uint64_t) n_embd || t->shape[1] != (uint64_t) n_vocab ||
            !strata::kernels::iq_supported((int) t->type) || n_embd % 256) {
            err = "native embedding: token_embd.weight is absent, of another shape, or of a type without a GPU "
@@ -92,11 +92,12 @@ bool expert_layout_load(const std::string& pack_dir, int64_t n_layers, int64_t n
    L.bytes.assign((size_t) n_layers, 0);
    L.max_blob = 0;
    std::string line;
+    // `layer gu_type d_type offset blob_bytes`; older packs' extra columns (shard-1 offsets) are not read
    while (std::getline(in, line)) {
        if (line.empty() || line[0] == '#') continue;
        std::istringstream ss(line);
        long long l = -1, gt = -1, dt = -1;
-        unsigned long long off = 0, blob = 0, go = 0, uo = 0, dox = 0;
+        unsigned long long off = 0, blob = 0;
        if (!(ss >> l >> gt >> dt >> off >> blob) || l < 0 || l >= n_layers) {
            err = "native_experts.txt: a malformed line: " + line;
            return false;
@@ -108,12 +109,6 @@ bool expert_layout_load(const std::string& pack_dir, int64_t n_layers, int64_t n
                  " B but its formats make " + std::to_string(f.bytes);
            return false;
        }
-        if (ss >> go >> uo >> dox) {   // v2 lines: the GGUF offsets
-            if (L.gguf_off.empty()) L.gguf_off.assign((size_t) (3 * n_layers), 0);
-            L.gguf_off[(size_t) (3 * l)] = go;
-            L.gguf_off[(size_t) (3 * l + 1)] = uo;
-            L.gguf_off[(size_t) (3 * l + 2)] = dox;
-        }
        L.fmt[(size_t) l] = f;
        L.offset[(size_t) l] = off;
        L.bytes[(size_t) l] = blob;
@@ -268,6 +268,228 @@ __device__ __forceinline__ float vec_dot_iq4_nl_q8_1(const void* __restrict__ vb
    return d * sumi;
 }

+// Q4_K / Q5_K gate/up and Q5_1 / Q8_0 down (the unsloth UD-Q4_K_XL file's experts)
+__device__ __forceinline__ float vec_dot_q4_K_q8_1_impl_vmmq(const int* __restrict__ v, const int* __restrict__ u,
+                                                             const uint8_t* __restrict__ sc, const uint8_t* __restrict__ m,
+                                                             const half2& dm4, const float* __restrict__ d8) {
+    float sumf_d = 0.0f;
+    float sumf_m = 0.0f;
+#pragma unroll
+    for (int i = 0; i < QR4_K; ++i) {
+        const int v0i = (v[0] >> (4 * i)) & 0x0F0F0F0F;
+        const int v1i = (v[1] >> (4 * i)) & 0x0F0F0F0F;
+        const int dot1 = ggml_cuda_dp4a(v1i, u[2 * i + 1], ggml_cuda_dp4a(v0i, u[2 * i + 0], 0));
+        const int dot2 = ggml_cuda_dp4a(0x01010101, u[2 * i + 1], ggml_cuda_dp4a(0x01010101, u[2 * i + 0], 0));
+        sumf_d += d8[i] * (dot1 * sc[i]);
+        sumf_m += d8[i] * (dot2 * m[i]);
+    }
+    const float2 dm4f = __half22float2(dm4);
+    return dm4f.x * sumf_d - dm4f.y * sumf_m;
+}
+
+// the 6-bit scale and min of the 32-value group pair `bq8_offset / 2` (vecdotq.cuh, shared by Q4_K and Q5_K)
+__device__ __forceinline__ void k_scales(const uint8_t* scales8, int bq8_offset, uint16_t aux[2]) {
+    const uint16_t* scales = (const uint16_t*) scales8;
+    const int j = bq8_offset / 2;
+    const int jm = j & 1;
+    const uint32_t s0 = scales[jm + 0];
+    const uint32_t s2 = scales[jm + 2];
+    const uint32_t s4 = scales[jm + 4];
+    const uint32_t hi = (uint32_t) -(int32_t) (j >= 2);
+    aux[0] = (uint16_t) (((s0 & 0x3f3f) & ~hi) | ((((s4 >> 0) & 0x0f0f) | ((s0 & 0xc0c0) >> 2)) & hi));
+    aux[1] = (uint16_t) (((s2 & 0x3f3f) & ~hi) | ((((s4 >> 4) & 0x0f0f) | ((s2 & 0xc0c0) >> 2)) & hi));
+}
+
+__device__ __forceinline__ float vec_dot_q4_K_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ bq8_1,
+                                                   const int& kbx, const int& iqs) {
+    const block_q4_K* bq4_K = (const block_q4_K*) vbq + kbx;
+    int v[2];
+    int u[2 * QR4_K];
+    float d8[QR4_K];
+    const int bq8_offset = QR4_K * ((iqs / 2) / (QI8_1 / 2));
+    const int* q4 = (const int*) (bq4_K->qs + 16 * bq8_offset + 4 * ((iqs / 2) % 4));
+    v[0] = q4[0];
+    v[1] = q4[4];
+    uint16_t aux[2];
+    k_scales(bq4_K->scales, bq8_offset, aux);
+    const uint8_t* sc = (const uint8_t*) aux;
+    const uint8_t* m = sc + 2;
+    for (int i = 0; i < QR4_K; ++i) {
+        const block_q8_1* bq8i = bq8_1 + bq8_offset + i;
+        d8[i] = __low2float(bq8i->ds);
+        const int* q8 = (const int*) bq8i->qs + ((iqs / 2) % 4);
+        u[2 * i + 0] = q8[0];
+        u[2 * i + 1] = q8[4];
+    }
+    return vec_dot_q4_K_q8_1_impl_vmmq(v, u, sc, m, bq4_K->dm, d8);
+}
+
+__device__ __forceinline__ float vec_dot_q5_K_q8_1_impl_vmmq(const int* __restrict__ vl, const int* __restrict__ vh,
+                                                             const int* __restrict__ u, const uint8_t* __restrict__ sc,
+                                                             const uint8_t* __restrict__ m, const half2& dm5,
+                                                             const float* __restrict__ d8) {
+    float sumf_d = 0.0f;
+    float sumf_m = 0.0f;
+#pragma unroll
+    for (int i = 0; i < QR5_K; ++i) {
+        const int vl0i = (vl[0] >> (4 * i)) & 0x0F0F0F0F;
+        const int vl1i = (vl[1] >> (4 * i)) & 0x0F0F0F0F;
+        const int vh0i = ((vh[0] >> i) << 4) & 0x10101010;
+        const int vh1i = ((vh[1] >> i) << 4) & 0x10101010;
+        const int v0i = vl0i | vh0i;
+        const int v1i = vl1i | vh1i;
+        const int dot1 = ggml_cuda_dp4a(v0i, u[2 * i + 0], ggml_cuda_dp4a(v1i, u[2 * i + 1], 0));
+        const int dot2 = ggml_cuda_dp4a(0x01010101, u[2 * i + 0], ggml_cuda_dp4a(0x01010101, u[2 * i + 1], 0));
+        sumf_d += d8[i] * (dot1 * sc[i]);
+        sumf_m += d8[i] * (dot2 * m[i]);
+    }
+    const float2 dm5f = __half22float2(dm5);
+    return dm5f.x * sumf_d - dm5f.y * sumf_m;
+}
+
+__device__ __forceinline__ float vec_dot_q5_K_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ bq8_1,
+                                                   const int& kbx, const int& iqs) {
+    const block_q5_K* bq5_K = (const block_q5_K*) vbq + kbx;
+    int vl[2];
+    int vh[2];
+    int u[2 * QR5_K];
+    float d8[QR5_K];
+    const int bq8_offset = QR5_K * ((iqs / 2) / (QI8_1 / 2));
+    const int* ql = (const int*) (bq5_K->qs + 16 * bq8_offset + 4 * ((iqs / 2) % 4));
+    const int* qh = (const int*) (bq5_K->qh + 4 * ((iqs / 2) % 4));
+    vl[0] = ql[0];
+    vl[1] = ql[4];
+    vh[0] = qh[0] >> bq8_offset;
+    vh[1] = qh[4] >> bq8_offset;
+    uint16_t aux[2];
+    k_scales(bq5_K->scales, bq8_offset, aux);
+    const uint8_t* sc = (const uint8_t*) aux;
+    const uint8_t* m = sc + 2;
+#pragma unroll
+    for (int i = 0; i < QR5_K; ++i) {
+        const block_q8_1* bq8i = bq8_1 + bq8_offset + i;
+        d8[i] = __low2float(bq8i->ds);
+        const int* q8 = (const int*) bq8i->qs + ((iqs / 2) % 4);
+        u[2 * i + 0] = q8[0];
+        u[2 * i + 1] = q8[4];
+    }
+    return vec_dot_q5_K_q8_1_impl_vmmq(vl, vh, u, sc, m, bq5_K->dm, d8);
+}
+
+constexpr int VDR_Q5_1 = 2, VDR_Q8_0 = 2;   // VDR_Q5_1_Q8_1_MMVQ, VDR_Q8_0_Q8_1_MMVQ
+
+// The weights are llama.cpp's; the min term is not.  llama.cpp multiplies the Q5_1 min by the q8_1 block's sum of
+// the ORIGINAL activations while the scaled term uses the quantized ones, which adds min * (sum x - sum x_q) to
+// every block: 1.9% instead of 1.2% relative error per expert (native_expert_parity).  Like ggml-cpu, and like
+// llama.cpp's own Q4_K / Q5_K dots, the min multiplies the quantized sum here.
+__device__ __forceinline__ float vec_dot_q5_1_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ bq8_1,
+                                                   const int& kbx, const int& iqs) {
+    const block_q5_1* bq5_1 = (const block_q5_1*) vbq + kbx;
+    int sumi = 0, sumu = 0;
+#pragma unroll
+    for (int i = 0; i < VDR_Q5_1; ++i) {
+        const int vl = get_int_b4(bq5_1->qs, iqs + i);
+        const int vh = get_int_b4(bq5_1->qh, 0) >> (4 * (iqs + i));
+        const int u0 = get_int_b4(bq8_1->qs, iqs + i), u1 = get_int_b4(bq8_1->qs, iqs + i + QI5_1);
+        int vi0 = (vl >> 0) & 0x0F0F0F0F;
+        vi0 |= (vh << 4) & 0x00000010;
+        vi0 |= (vh << 11) & 0x00001000;
+        vi0 |= (vh << 18) & 0x00100000;
+        vi0 |= (vh << 25) & 0x10000000;
+        sumi = ggml_cuda_dp4a(vi0, u0, sumi);
+        int vi1 = (vl >> 4) & 0x0F0F0F0F;
+        vi1 |= (vh >> 12) & 0x00000010;
+        vi1 |= (vh >> 5) & 0x00001000;
+        vi1 |= (vh << 2) & 0x00100000;
+        vi1 |= (vh << 9) & 0x10000000;
+        sumi = ggml_cuda_dp4a(vi1, u1, sumi);
+        sumu = ggml_cuda_dp4a(0x01010101, u1, ggml_cuda_dp4a(0x01010101, u0, sumu));
+    }
+    const float2 dm5 = __half22float2(bq5_1->dm);
+    const float d8 = __low2float(bq8_1->ds);
+    return sumi * (dm5.x * d8) + sumu * (dm5.y * d8);
+}
+
+// Q6_K gate/up (UD-Q5_K_XL, UD-Q6_K_XL), VDR_Q6_K_Q8_1_MMVQ = 1
+__device__ __forceinline__ float vec_dot_q6_K_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ bq8_1,
+                                                   const int& kbx, const int& iqs) {
+    const block_q6_K* bq6_K = (const block_q6_K*) vbq + kbx;
+    const int bq8_offset = 2 * QR6_K * (iqs / (QI6_K / 2)) + (iqs % (QI6_K / 2)) / (QI6_K / 4);
+    const int scale_offset = (QI6_K / 4) * (iqs / (QI6_K / 2)) + (iqs % (QI6_K / 2)) / (QI6_K / 8);
+    const int vh_shift = 2 * ((iqs % (QI6_K / 2)) / (QI6_K / 4));
+    const int vl = get_int_b2(bq6_K->ql, iqs);
+    const int vh = get_int_b2(bq6_K->qh, (QI6_K / 4) * (iqs / (QI6_K / 2)) + iqs % (QI6_K / 4)) >> vh_shift;
+    const int8_t* scales = bq6_K->scales + scale_offset;
+    float sumf = 0.0f;
+#pragma unroll
+    for (int i = 0; i < QR6_K; ++i) {
+        const int u = get_int_b4(bq8_1[bq8_offset + 2 * i].qs, iqs % QI8_1);
+        const float d8 = __low2float(bq8_1[bq8_offset + 2 * i].ds);
+        const int sc = scales[4 * i];
+        const int vil = (vl >> (4 * i)) & 0x0F0F0F0F;
+        const int vih = ((vh >> (4 * i)) << 4) & 0x30303030;
+        const int vi = __vsubss4((vil | vih), 0x20202020);   // (vil | vih) - 32
+        sumf += d8 * (ggml_cuda_dp4a(vi, u, 0) * sc);
+    }
+    return __half2float(bq6_K->d) * sumf;
+}
+
+// IQ1_S gate/up (UD-IQ1_S), VDR 1.  As for Q5_1, the delta term takes the quantized activation sum (ggml-cpu's
+// bsums) rather than llama.cpp's sum of the original values.
+__device__ __forceinline__ float vec_dot_iq1_s_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ bq8_1,
+                                                    const int& kbx, const int& iqs) {
+    const block_iq1_s* bq1 = (const block_iq1_s*) vbq + kbx;
+    const int qs_packed = get_int_b2(bq1->qs, iqs);
+    const uint8_t* qs = (const uint8_t*) &qs_packed;
+    const int qh = bq1->qh[iqs];
+    int sumi = 0, sumy = 0;
+#pragma unroll
+    for (int l0 = 0; l0 < 8; l0 += 2) {
+        const int grid = iq1s_grid_gpu[qs[l0 / 2] | (((qh >> 3 * (l0 / 2)) & 0x07) << 8)];
+        const int grid0 = (grid >> 0) & 0x0F0F0F0F;
+        const int grid1 = (grid >> 4) & 0x0F0F0F0F;
+        const int u0 = get_int_b4(bq8_1[iqs].qs, l0 + 0);
+        const int u1 = get_int_b4(bq8_1[iqs].qs, l0 + 1);
+        sumi = ggml_cuda_dp4a(grid0, u0, sumi);
+        sumi = ggml_cuda_dp4a(grid1, u1, sumi);
+        sumy = ggml_cuda_dp4a(0x01010101, u1, ggml_cuda_dp4a(0x01010101, u0, sumy));
+    }
+    const float d1q = __half2float(bq1->d) * (((qh >> 11) & 0x0E) + 1);
+    const float delta = -1.0f + IQ1S_DELTA - (qh & 0x8000) * (2.0f * IQ1S_DELTA / 0x8000);
+    return d1q * __low2float(bq8_1[iqs].ds) * (sumi + sumy * delta);
+}
+
+// IQ4_XS gate/up (UD-IQ4_XS, UD-Q3_K_XL), VDR 4
+__device__ __forceinline__ float vec_dot_iq4_xs_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ bq8_1,
+                                                     const int& kbx, const int& iqs) {
+    const block_iq4_xs* bq4 = (const block_iq4_xs*) vbq + kbx;
+    int sumi = 0;
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
+        const int aux_q4 = get_int_b4(bq4->qs, iqs + j);
+        const int2 v = get_int_from_table_16(aux_q4, kvalues_iq4nl);
+        const int u0 = get_int_b4(bq8_1[iqs / 4].qs, j + 0);
+        const int u1 = get_int_b4(bq8_1[iqs / 4].qs, j + 4);
+        sumi = ggml_cuda_dp4a(v.x, u0, sumi);
+        sumi = ggml_cuda_dp4a(v.y, u1, sumi);
+    }
+    const int ls = ((bq4->scales_l[iqs / 8] >> (iqs & 0x04)) & 0x0F) | (((bq4->scales_h >> (iqs / 2)) & 0x03) << 4);
+    sumi *= ls - 32;
+    const float d = __half2float(bq4->d) * __low2float(bq8_1[iqs / 4].ds);
+    return d * sumi;
+}
+
+__device__ __forceinline__ float vec_dot_q8_0_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ bq8_1,
+                                                   const int& kbx, const int& iqs) {
+    const block_q8_0* bq8_0 = (const block_q8_0*) vbq + kbx;
+    int sumi = 0;
+#pragma unroll
+    for (int i = 0; i < VDR_Q8_0; ++i)
+        sumi = ggml_cuda_dp4a(get_int_b2(bq8_0->qs, iqs + i), get_int_b4(bq8_1->qs, iqs + i), sumi);
+    const float d8_0 = __half2float(bq8_0->d), d8_1 = __low2float(bq8_1->ds);
+    return d8_0 * d8_1 * ((float) sumi);
+}
+
 // ---------------------------------------------------------------- the formats
 // qk = values per block, ipb = dot calls per block (qi / vdr), step = the iqs stride between calls.
 template<int TY> struct Fmt;
@@ -287,6 +509,32 @@ template<> struct Fmt<29> { static constexpr int qk = 256, ipb = 8, step = 1;
    __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_iq1_m_q8_1(v, y, kbx, iqs); } };
 template<> struct Fmt<42> { static constexpr int qk = 64, ipb = 2, step = 1;
    __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_q2_0_q8_1(v, y, kbx, iqs); } };
+template<> struct Fmt<12> { static constexpr int qk = 256, ipb = QI4_K / 2, step = 2;
+    __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_q4_K_q8_1(v, y, kbx, iqs); } };
+template<> struct Fmt<13> { static constexpr int qk = 256, ipb = QI5_K / 2, step = 2;
+    __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_q5_K_q8_1(v, y, kbx, iqs); } };
+template<> struct Fmt<7> { static constexpr int qk = 32, ipb = QI5_1 / VDR_Q5_1, step = VDR_Q5_1;
+    __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_q5_1_q8_1(v, y, kbx, iqs); } };
+template<> struct Fmt<8> { static constexpr int qk = 32, ipb = QI8_0 / VDR_Q8_0, step = VDR_Q8_0;
+    __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_q8_0_q8_1(v, y, kbx, iqs); } };
+template<> struct Fmt<14> { static constexpr int qk = 256, ipb = QI6_K, step = 1;
+    __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_q6_K_q8_1(v, y, kbx, iqs); } };
+template<> struct Fmt<19> { static constexpr int qk = 256, ipb = QI1_S, step = 1;
+    __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_iq1_s_q8_1(v, y, kbx, iqs); } };
+template<> struct Fmt<23> { static constexpr int qk = 256, ipb = QI4_XS / 4, step = 4;
+    __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_iq4_xs_q8_1(v, y, kbx, iqs); } };
+
+// Every format above, for the dispatch switches: X(type) once per Fmt<type>.
+#define STRATA_FMTS(X) X(7) X(8) X(12) X(13) X(14) X(16) X(17) X(18) X(19) X(20) X(21) X(22) X(23) X(29) X(42)
+
+int fmt_qk(int t) {
+    switch (t) {
+#define STRATA_QK(T) case T: return Fmt<T>::qk;
+        STRATA_FMTS(STRATA_QK)
+#undef STRATA_QK
+        default: return 0;
+    }
+}

 __device__ __forceinline__ float warp_sum(float v) {
 #pragma unroll
@@ -543,6 +791,113 @@ __device__ void dq_q2_0(const void* vx, int64_t ibs, dst_t* yy, int tid) {
        yy[b * 64 + i] = cvt<dst_t>(d * (float) (code - 1));
    }
 }
+__device__ __forceinline__ void scale_min_k4(int j, const uint8_t* q, uint8_t& d, uint8_t& m) {
+    if (j < 4) {
+        d = q[j] & 63; m = q[j + 4] & 63;
+    } else {
+        d = (q[j + 4] & 0xF) | ((q[j - 4] >> 6) << 4);
+        m = (q[j + 4] >> 4) | ((q[j - 0] >> 6) << 4);
+    }
+}
+template<typename dst_t>
+__device__ void dq_q4_k(const void* vx, int64_t ibs, dst_t* yy, int tid) {
+    const block_q4_K* x = (const block_q4_K*) vx;
+    const int64_t il = tid / 8, ir = tid % 8, is = 2 * il;
+    const int n = 4;
+    dst_t* y = yy + 64 * il + n * ir;
+    const float dall = __low2half(x[ibs].dm);
+    const float dmin = __high2half(x[ibs].dm);
+    const uint8_t* q = x[ibs].qs + 32 * il + n * ir;
+    uint8_t sc, m;
+    scale_min_k4((int) is + 0, x[ibs].scales, sc, m);
+    const float d1 = dall * sc, m1 = dmin * m;
+    scale_min_k4((int) is + 1, x[ibs].scales, sc, m);
+    const float d2 = dall * sc, m2 = dmin * m;
+    for (int l = 0; l < n; ++l) {
+        y[l + 0] = cvt<dst_t>(d1 * (q[l] & 0xF) - m1);
+        y[l + 32] = cvt<dst_t>(d2 * (q[l] >> 4) - m2);
+    }
+}
+// llama.cpp's dequantize_q5_K, its 64 threads folded onto 32
+template<typename dst_t>
+__device__ void dq_q5_k(const void* vx, int64_t ibs, dst_t* yy, int tid) {
+    const block_q5_K* x = (const block_q5_K*) vx;
+    for (int tt = tid; tt < 64; tt += 32) {
+        const int il = tt / 16, ir = tt % 16, is = 2 * il;
+        dst_t* y = yy + 64 * il + 2 * ir;
+        const float dall = __low2half(x[ibs].dm);
+        const float dmin = __high2half(x[ibs].dm);
+        const uint8_t* ql = x[ibs].qs + 32 * il + 2 * ir;
+        const uint8_t* qh = x[ibs].qh + 2 * ir;
+        uint8_t sc, m;
+        scale_min_k4(is + 0, x[ibs].scales, sc, m);
+        const float d1 = dall * sc, m1 = dmin * m;
+        scale_min_k4(is + 1, x[ibs].scales, sc, m);
+        const float d2 = dall * sc, m2 = dmin * m;
+        uint8_t hm = (uint8_t) (1 << (2 * il));
+        y[0] = cvt<dst_t>(d1 * ((ql[0] & 0xF) + (qh[0] & hm ? 16 : 0)) - m1);
+        y[1] = cvt<dst_t>(d1 * ((ql[1] & 0xF) + (qh[1] & hm ? 16 : 0)) - m1);
+        hm <<= 1;
+        y[32] = cvt<dst_t>(d2 * ((ql[0] >> 4) + (qh[0] & hm ? 16 : 0)) - m2);
+        y[33] = cvt<dst_t>(d2 * ((ql[1] >> 4) + (qh[1] & hm ? 16 : 0)) - m2);
+    }
+}
+// 32-value blocks: a "superblock" is 8 of them; thread tid writes 8 values of block tid % 8
+template<typename dst_t>
+__device__ void dq_q5_1(const void* vx, int64_t ibs, dst_t* yy, int tid) {
+    const block_q5_1* x = (const block_q5_1*) vx + ibs * (QK_K / QK5_1);
+    const int ib = tid % 8, il = tid / 8;
+    const float2 dm = __half22float2(x[ib].dm);
+    uint32_t qh;
+    memcpy(&qh, x[ib].qh, sizeof(qh));
+    dst_t* y = yy + 32 * ib;
+    for (int j = 0; j < 4; ++j) {
+        const int iqs = 4 * il + j;
+        const int xh_0 = ((qh >> (iqs + 0)) << 4) & 0x10;
+        const int xh_1 = ((qh >> (iqs + 12))) & 0x10;
+        y[iqs] = cvt<dst_t>((float) ((x[ib].qs[iqs] & 0xf) | xh_0) * dm.x + dm.y);
+        y[iqs + 16] = cvt<dst_t>((float) ((x[ib].qs[iqs] >> 4) | xh_1) * dm.x + dm.y);
+    }
+}
+template<typename dst_t>
+__device__ void dq_q8_0(const void* vx, int64_t ibs, dst_t* yy, int tid) {
+    const block_q8_0* x = (const block_q8_0*) vx + ibs * (QK_K / QK8_0);
+    const int ib = tid % 8, il = tid / 8;
+    const float d = __half2float(x[ib].d);
+    dst_t* y = yy + 32 * ib + 8 * il;
+    for (int j = 0; j < 8; ++j) y[j] = cvt<dst_t>((float) x[ib].qs[8 * il + j] * d);
+}
+// llama.cpp's dequantize_q6_K, its 64 threads folded onto 32
+template<typename dst_t>
+__device__ void dq_q6_k(const void* vx, int64_t ibs, dst_t* yy, int tid) {
+    const block_q6_K* x = (const block_q6_K*) vx;
+    for (int tt = tid; tt < 64; tt += 32) {
+        const int ip = tt / 32, il = tt - 32 * ip, is = 8 * ip + il / 16;
+        dst_t* y = yy + 128 * ip + il;
+        const float d = __half2float(x[ibs].d);
+        const uint8_t* ql = x[ibs].ql + 64 * ip + il;
+        const uint8_t qh = x[ibs].qh[32 * ip + il];
+        const int8_t* sc = x[ibs].scales + is;
+        y[0] = cvt<dst_t>(d * sc[0] * ((int8_t) ((ql[0] & 0xF) | (((qh >> 0) & 3) << 4)) - 32));
+        y[32] = cvt<dst_t>(d * sc[2] * ((int8_t) ((ql[32] & 0xF) | (((qh >> 2) & 3) << 4)) - 32));
+        y[64] = cvt<dst_t>(d * sc[4] * ((int8_t) ((ql[0] >> 4) | (((qh >> 4) & 3) << 4)) - 32));
+        y[96] = cvt<dst_t>(d * sc[6] * ((int8_t) ((ql[32] >> 4) | (((qh >> 6) & 3) << 4)) - 32));
+    }
+}
+template<typename dst_t>
+__device__ void dq_iq1_s(const void* vx, int64_t ibs, dst_t* yy, int tid) {
+    const block_iq1_s* x = (const block_iq1_s*) vx;
+    const int64_t il = tid / 8, ib = tid % 8;
+    dst_t* y = yy + 32 * ib + 8 * il;
+    const float delta = x[ibs].qh[ib] & 0x8000 ? -1 - IQ1S_DELTA : -1 + IQ1S_DELTA;
+    const float d = (float) x[ibs].d * (2 * ((x[ibs].qh[ib] >> 12) & 7) + 1);
+    uint32_t grid32[2];
+    const int8_t* q = (const int8_t*) grid32;
+    grid32[0] = iq1s_grid_gpu[x[ibs].qs[4 * ib + il] | (((x[ibs].qh[ib] >> 3 * il) & 7) << 8)];
+    grid32[1] = (grid32[0] >> 4) & 0x0f0f0f0f;
+    grid32[0] &= 0x0f0f0f0f;
+    for (int j = 0; j < 8; ++j) y[j] = cvt<dst_t>(d * (q[j] + delta));
+}

 template<typename dst_t>
 __device__ __forceinline__ void dq_dispatch(int ty, const void* vx, int64_t ibs, dst_t* y, int tid) {
@@ -557,6 +912,12 @@ __device__ __forceinline__ void dq_dispatch(int ty, const void* vx, int64_t ibs,
        case 23: dq_iq4_xs(vx, ibs, y, tid); break;
        case 11: dq_q3_k(vx, ibs, y, tid); break;
        case 42: dq_q2_0(vx, ibs, y, tid); break;
+        case 12: dq_q4_k(vx, ibs, y, tid); break;
+        case 13: dq_q5_k(vx, ibs, y, tid); break;
+        case 7: dq_q5_1(vx, ibs, y, tid); break;
+        case 8: dq_q8_0(vx, ibs, y, tid); break;
+        case 14: dq_q6_k(vx, ibs, y, tid); break;
+        case 19: dq_iq1_s(vx, ibs, y, tid); break;
        default: break;
    }
 }
@@ -576,7 +937,11 @@ __global__ void dequant_gu_kernel(int ty, const void* __restrict__ gate, const v
    dq_dispatch<__half>(ty, parity ? up : gate, i, y + ((2 * r + parity) * per_row + c) * QK_K, threadIdx.x);
 }

-bool is_iq(int t) { return t == 16 || t == 17 || t == 18 || t == 20 || t == 21 || t == 22 || t == 23 || t == 29 || t == 42 || t == 11; }
+// the formats dq_dispatch dequantizes
+bool is_iq(int t) {
+    return t == 16 || t == 17 || t == 18 || t == 20 || t == 21 || t == 22 || t == 23 || t == 29 || t == 42 || t == 11 ||
+           t == 12 || t == 13 || t == 14 || t == 19 || t == 7 || t == 8;
+}

 }  // namespace

@@ -594,6 +959,12 @@ size_t iq_row_bytes(int t, int64_t n) noexcept {
        case 23: return (size_t) (n / 256) * sizeof(block_iq4_xs);
        case 11: return (size_t) (n / 256) * sizeof(block_q3_K);
        case 42: return (size_t) (n / 64) * sizeof(block_q2_0);
+        case 12: return (size_t) (n / 256) * sizeof(block_q4_K);
+        case 13: return (size_t) (n / 256) * sizeof(block_q5_K);
+        case 14: return (size_t) (n / 256) * sizeof(block_q6_K);
+        case 19: return (size_t) (n / 256) * sizeof(block_iq1_s);
+        case 7: return (size_t) (n / 32) * sizeof(block_q5_1);
+        case 8: return (size_t) (n / 32) * sizeof(block_q8_0);
        default: return 0;
    }
 }
@@ -612,14 +983,9 @@ void iq_mmvq(int t, const void* w, const void* x_q8_1, float* y, int n_in, int n
    const auto* W = (const uint8_t*) w;
    const auto* X = (const block_q8_1*) x_q8_1;
    switch (t) {
-        case 16: mmvq_kernel<16><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
-        case 17: mmvq_kernel<17><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
-        case 18: mmvq_kernel<18><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
-        case 20: mmvq_kernel<20><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
-        case 21: mmvq_kernel<21><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
-        case 22: mmvq_kernel<22><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
-        case 29: mmvq_kernel<29><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
-        case 42: mmvq_kernel<42><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
+#define STRATA_MMVQ(T) case T: mmvq_kernel<T><<<grid, block, 0, s>>>(W, rb, X, y, n_in, n_out, ncols); break;
+        STRATA_FMTS(STRATA_MMVQ)
+#undef STRATA_MMVQ
        default: std::fprintf(stderr, "iq_mmvq: type %d is not supported\n", t); std::exit(1);
    }
    check("iq_mmvq");
@@ -663,6 +1029,11 @@ void iq_dequant_gu_f16(int t, const void* gate, const void* up, int64_t n_ff, in
    check("iq_dequant_gu_f16");
 }

+bool native_expert_supported(int gu_type, int d_type, int64_t n_embd, int64_t n_ff) noexcept {
+    const int qg = fmt_qk(gu_type), qd = fmt_qk(d_type);
+    return qg > 0 && qd > 0 && is_iq(gu_type) && is_iq(d_type) && n_embd % qg == 0 && n_ff % qd == 0;
+}
+
 NativeExpertLayout native_expert_layout(int gu_type, int d_type, int64_t n_embd, int64_t n_ff) {
    NativeExpertLayout L;
    L.gu_type = gu_type;
@@ -695,13 +1066,9 @@ void native_expert_grouped(const NativeExpertLayout& L, const unsigned long long
    const auto* X = (const block_q8_1*) x_q8_1;
    const dim3 ggu((unsigned) ((2 * L.n_ff + GU_ROWS - 1) / GU_ROWS), (unsigned) cap_groups);
    switch (L.gu_type) {
-        case 16: native_gu_kernel<16><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
-        case 17: native_gu_kernel<17><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
-        case 18: native_gu_kernel<18><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
-        case 21: native_gu_kernel<21><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
-        case 22: native_gu_kernel<22><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
-        case 29: native_gu_kernel<29><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
-        case 42: native_gu_kernel<42><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
+#define STRATA_GU(T) case T: native_gu_kernel<T><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
+        STRATA_FMTS(STRATA_GU)
+#undef STRATA_GU
        default: std::fprintf(stderr, "native_expert_grouped: gate/up type %d\n", L.gu_type); std::exit(1);
    }
    check("native_expert_grouped/gu");
@@ -710,8 +1077,9 @@ void native_expert_grouped(const NativeExpertLayout& L, const unsigned long long
    quantize_q8_1_kernel<<<(unsigned) ((nh + 255) / 256), 256, 0, s>>>(h, hq, nh);
    const dim3 gd((unsigned) ((L.n_embd + 7) / 8), (unsigned) cap_groups);
    switch (L.d_type) {
-        case 20: native_down_kernel<20><<<gd, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); break;
-        case 42: native_down_kernel<42><<<gd, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); break;
+#define STRATA_DOWN(T) case T: native_down_kernel<T><<<gd, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); break;
+        STRATA_FMTS(STRATA_DOWN)
+#undef STRATA_DOWN
        default: std::fprintf(stderr, "native_expert_grouped: down type %d\n", L.d_type); std::exit(1);
    }
    check("native_expert_grouped/down");
@@ -228,8 +228,8 @@ void ple_block(const float* emb, const float* hidden, const float* hist_rows, co
               PleOut& out, void* scratch, void* stream) {
    const bool native_key = w.key_native_data != nullptr && w.key_bf16 == nullptr;
    if (native_key && (!emb || !hidden || !hist_rows || !out.result || !scratch || !stream ||
-                       !w.key_native_q8_1 || w.key_native_type != 42))
-        throw std::invalid_argument("ple_block: native key requires Q2_0 weights, input/output, private scratch and explicit stream");
+                       !w.key_native_q8_1 || !native_mmvq_supported(w.key_native_type)))
+        throw std::invalid_argument("ple_block: native key requires native MMVQ weights, input/output, private scratch and explicit stream");
    if (emb == nullptr || hidden == nullptr || hist_rows == nullptr || out.result == nullptr) return;
    const int n_embd = NG_N_EMBD, hc = NG_HC, hc_dim = NG_HC_DIM;
    static_assert(NG_N_EMBD == 2560 && NG_HC_DIM == 10240, "native PLE key geometry changed");
@@ -1,10 +1,12 @@
-// src/kernels/native_expert_parity.cpp - plan v0.3 P6: one native expert three ways, on real GGUF rows.
+// src/kernels/native_expert_parity.cpp - plan v0.3 P6: one native expert three ways.
 //
-//     build/native_expert_parity <shard1.gguf> [layer ...]
+//     build/native_expert_parity <shard1.gguf> [layer ...]      real rows (a split model's shards are found by name)
+//     build/native_expert_parity --synthetic GU/DOWN ...        random weights quantized by ggml, e.g. iq1_s/iq4_nl
 //
 // (a) float reference: ggml's own dequantizer (`to_float`) and a float SwiGLU expert, (b) the CPU path
 // (ggml-cpu vec_dot with its quantized activations), (c) the GPU path (`native_expert_grouped`, q8_1
-// activations).  (b) and (c) each differ from (a) by their activation rounding only (a few 1e-3 relative).
+// activations).  (b) and (c) each differ from (a) by their activation rounding only (about 1e-2 relative).
+// (d) the GPU dequantizers of the prompt path and the embedding, against `to_float`: exact.
 #include "strata/artifact/gguf_reader.hpp"
 #include "strata/kernels/cpu/native_expert.hpp"
 #include "strata/kernels/cpu/expert.hpp"
@@ -16,6 +18,7 @@

 #include <cuda_runtime.h>

+#include <cctype>
 #include <chrono>
 #include <cmath>
 #include <cstdio>
@@ -26,178 +29,266 @@

 namespace cpu = strata::kernels::cpu;

-static double rel(const std::vector<float>& a, const std::vector<float>& b) {
+namespace {
+
+constexpr int NT = 3;
+constexpr int64_t H = 2560, FF = 640;
+
+double rel(const std::vector<float>& a, const std::vector<float>& b) {
    double n = 0, d = 0;
    for (size_t i = 0; i < a.size(); ++i) { n += std::fabs((double) a[i] - b[i]); d += std::fabs((double) b[i]); }
    return n / (d + 1e-30);
 }

+double us_since(std::chrono::steady_clock::time_point t0, int it) {
+    return std::chrono::duration<double, std::micro>(std::chrono::steady_clock::now() - t0).count() / it;
+}
+
+// The three ways and the dequantizers for one expert blob; returns the number of failed checks.
+int check_blob(const cpu::NativeFmt& f, const std::vector<uint8_t>& blob, int seed, const std::string& label,
+               cudaStream_t s) {
+    int failures = 0;
+    // (a) the float reference
+    const auto* tg = ggml_get_type_traits((ggml_type) f.gu_type);
+    const auto* td = ggml_get_type_traits((ggml_type) f.d_type);
+    std::vector<float> G((size_t) FF * H), U((size_t) FF * H), D((size_t) H * FF);
+    for (int64_t r = 0; r < FF; ++r) {
+        tg->to_float(blob.data() + r * f.gu_row, G.data() + r * H, H);
+        tg->to_float(blob.data() + f.up_off + r * f.gu_row, U.data() + r * H, H);
+    }
+    for (int64_t r = 0; r < H; ++r) td->to_float(blob.data() + f.down_off + r * f.d_row, D.data() + r * FF, FF);
+    std::mt19937 rng(11 + seed);
+    std::normal_distribution<float> nd(0.f, 1.f);
+    std::vector<float> x((size_t) NT * H);
+    for (auto& v : x) v = nd(rng);
+    std::vector<float> ref((size_t) NT * H), got_c((size_t) NT * H), got_g((size_t) NT * H);
+    for (int k = 0; k < NT; ++k) {
+        std::vector<float> h(FF);
+        for (int64_t r = 0; r < FF; ++r) {
+            double g = 0, u = 0;
+            for (int64_t i = 0; i < H; ++i) { g += (double) G[r * H + i] * x[k * H + i]; u += (double) U[r * H + i] * x[k * H + i]; }
+            h[r] = (float) (g / (1.0 + std::exp(-g)) * u);
+        }
+        for (int64_t r = 0; r < H; ++r) {
+            double o = 0;
+            for (int64_t i = 0; i < FF; ++i) o += (double) D[r * FF + i] * h[i];
+            ref[k * H + r] = (float) o;
+        }
+    }
+    // (b) the CPU
+    {
+        std::vector<std::vector<uint8_t>> act(NT, std::vector<uint8_t>(cpu::kNativeActBytes));
+        std::vector<std::vector<uint8_t>> hq(NT, std::vector<uint8_t>(cpu::kNativeHBytes));
+        std::vector<std::vector<float>> ff(NT, std::vector<float>(FF));
+        const void* a[NT];
+        float* ffp[NT];
+        const void* hp[NT];
+        float* op[NT];
+        for (int k = 0; k < NT; ++k) {
+            cpu::native_quant_act(f, x.data() + k * H, act[k].data());
+            a[k] = act[k].data();
+            ffp[k] = ff[k].data();
+        }
+        cpu::native_gu_rows(f, blob.data(), a, NT, ffp, 0, (int) FF);
+        const auto* tc = ggml_get_type_traits_cpu((ggml_type) f.gu_type);
+        const int it = 50;
+        if (cpu::iq512_supported(f.gu_type)) {
+            // the AVX-512 rows against ggml's own vec_dot, same Q8_K activations: float-order differences only
+            std::vector<float> g512((size_t) NT * FF), gref((size_t) NT * FF);
+            float* gp[NT];
+            for (int k = 0; k < NT; ++k) gp[k] = g512.data() + k * FF;
+            cpu::iq512_rows(f.gu_type, blob.data(), f.gu_row, (int) H, a, NT, gp, 0, (int) FF);
+            for (int k = 0; k < NT; ++k)
+                for (int64_t r = 0; r < FF; ++r)
+                    tc->vec_dot((int) H, &gref[k * FF + r], 0, blob.data() + r * f.gu_row, 0, a[k], 0, 1);
+            std::printf("          %s AVX-512 gate rows vs ggml vec_dot: rel %.2e\n",
+                        ggml_type_name((ggml_type) f.gu_type), rel(g512, gref));
+            float* f1[1] = {ffp[0]};
+            auto t0 = std::chrono::steady_clock::now();
+            for (int i = 0; i < it; ++i) cpu::iq512_gu_rows(f.gu_type, blob.data(), f.gu_row, f.up_off, (int) H, a, 1, f1, 0, (int) FF);
+            const double us1 = us_since(t0, it);
+            t0 = std::chrono::steady_clock::now();
+            for (int i = 0; i < it; ++i) cpu::iq512_gu_rows(f.gu_type, blob.data(), f.gu_row, f.up_off, (int) H, a, NT, ffp, 0, (int) FF);
+            std::printf("          gate+up one thread, weights cache-resident: avx512 %.0f us one token (%.2f GB/s), "
+                        "%.0f us %d tokens\n", us1, 2.0 * f.up_off / us1 / 1e3, us_since(t0, it), NT);
+        }
+        {   // ggml-cpu's own rows, one thread (what the pool runs for formats without an AVX-512 kernel)
+            auto t0 = std::chrono::steady_clock::now();
+            for (int i = 0; i < it; ++i)
+                for (int64_t r = 0; r < FF; ++r) {
+                    float gg, uu;
+                    tc->vec_dot((int) H, &gg, 0, blob.data() + r * f.gu_row, 0, a[0], 0, 1);
+                    tc->vec_dot((int) H, &uu, 0, blob.data() + f.up_off + r * f.gu_row, 0, a[0], 0, 1);
+                }
+            const double usg = us_since(t0, it);
+            std::printf("          gate+up one thread, weights cache-resident: ggml %.0f us one token (%.2f GB/s)\n", usg,
+                        2.0 * f.up_off / usg / 1e3);
+        }
+        for (int k = 0; k < NT; ++k) {
+            cpu::native_quant_h(f, ff[k].data(), hq[k].data());
+            hp[k] = hq[k].data();
+            op[k] = got_c.data() + k * H;
+        }
+        cpu::native_down_rows(f, blob.data(), hp, NT, op, 0, (int) H);
+        if (f.d_type == 42) {
+            // (b2) the AVX-512 GGUF-layout Q2_0 kernel the pool uses for Q2_0 down projections
+            std::vector<cpu::ActQ> a2(NT);
+            const cpu::ActQ* ap[NT];
+            std::vector<float> alt((size_t) NT * H);
+            float* altp[NT];
+            for (int k = 0; k < NT; ++k) {
+                cpu::act_quant_q8_1(ff[k].data(), (int) FF, a2[k]);
+                ap[k] = &a2[k];
+                altp[k] = alt.data() + k * H;
+            }
+            cpu::q2_0_gguf_rows_multi(blob.data() + f.down_off, f.d_row, (int) (FF / 64), ap, NT, altp, 0, (int) H);
+            std::printf("          q2_0 AVX-512 down vs ggml down: rel %.2e\n", rel(alt, got_c));
+        }
+    }
+    // (c) the GPU: one group holding the NT entries; (d) its dequantizers
+    {
+        const auto L = strata::kernels::native_expert_layout(f.gu_type, f.d_type, H, FF);
+        void *dblob, *dx, *dxq, *dscr;
+        float* dout;
+        unsigned long long* dptr;
+        int32_t *dstart, *dn, *ddst, *dtok;
+        cudaMalloc(&dblob, blob.size());
+        cudaMalloc(&dx, x.size() * 4);
+        cudaMalloc(&dxq, (size_t) NT * H / 32 * 36);
+        cudaMalloc(&dscr, strata::kernels::native_expert_scratch_bytes(NT, FF));
+        cudaMalloc((void**) &dout, (size_t) NT * H * 4);
+        cudaMalloc((void**) &dptr, 8);
+        cudaMalloc((void**) &dstart, 8);
+        cudaMalloc((void**) &dn, 4);
+        cudaMalloc((void**) &ddst, NT * 4);
+        cudaMalloc((void**) &dtok, NT * 4);
+        cudaMemcpy(dblob, blob.data(), blob.size(), cudaMemcpyHostToDevice);
+        cudaMemcpy(dx, x.data(), x.size() * 4, cudaMemcpyHostToDevice);
+        const unsigned long long p = (unsigned long long) dblob;
+        const int32_t st[2] = {0, NT}, one = 1, idx[NT] = {0, 1, 2};
+        cudaMemcpy(dptr, &p, 8, cudaMemcpyHostToDevice);
+        cudaMemcpy(dstart, st, 8, cudaMemcpyHostToDevice);
+        cudaMemcpy(dn, &one, 4, cudaMemcpyHostToDevice);
+        cudaMemcpy(ddst, idx, NT * 4, cudaMemcpyHostToDevice);
+        cudaMemcpy(dtok, idx, NT * 4, cudaMemcpyHostToDevice);
+        strata::kernels::quantize_q8_1_rows((const float*) dx, NT, H, dxq, s);
+        strata::kernels::native_expert_grouped(L, dptr, dstart, dn, ddst, dtok, 1, NT, dxq, dscr, dout, s);
+        cudaStreamSynchronize(s);
+        cudaMemcpy(got_g.data(), dout, got_g.size() * 4, cudaMemcpyDeviceToHost);
+        float* dq = nullptr;
+        cudaMalloc((void**) &dq, (size_t) FF * H * 4);
+        std::vector<float> got_dq((size_t) FF * H);
+        double dq_err = 0.0;
+        const struct { const uint8_t* src; int type; const std::vector<float>* want; } mats[3] = {
+            {(const uint8_t*) dblob, f.gu_type, &G}, {(const uint8_t*) dblob + f.up_off, f.gu_type, &U},
+            {(const uint8_t*) dblob + f.down_off, f.d_type, &D}};
+        for (const auto& m : mats) {
+            strata::kernels::iq_dequant_f32(m.type, m.src, FF * H, dq, s);
+            cudaStreamSynchronize(s);
+            cudaMemcpy(got_dq.data(), dq, got_dq.size() * 4, cudaMemcpyDeviceToHost);
+            dq_err = (std::max)(dq_err, rel(got_dq, *m.want));
+        }
+        std::printf("          GPU dequant vs ggml to_float: rel %.2e\n", dq_err);
+        if (!(dq_err < 1e-6)) ++failures;
+        cudaFree(dq);
+        cudaFree(dblob); cudaFree(dx); cudaFree(dxq); cudaFree(dscr); cudaFree(dout); cudaFree(dptr);
+        cudaFree(dstart); cudaFree(dn); cudaFree(ddst); cudaFree(dtok);
+    }
+    const double ec = rel(got_c, ref), eg = rel(got_g, ref), ecg = rel(got_c, got_g);
+    const bool ok = ec < 3e-2 && eg < 3e-2 && std::isfinite(ec) && std::isfinite(eg);
+    std::printf("%-9s %-8s/%-7s blob %8zu  cpu rel %.2e  gpu rel %.2e  cpu-gpu %.2e  %s\n", label.c_str(),
+                ggml_type_name((ggml_type) f.gu_type), ggml_type_name((ggml_type) f.d_type), f.bytes, ec, eg, ecg,
+                ok ? "ok" : "FAIL");
+    return failures + (ok ? 0 : 1);
+}
+
+int type_of(std::string name) {
+    for (auto& c : name) c = (char) std::tolower((unsigned char) c);
+    for (int t = 0; t < GGML_TYPE_COUNT; ++t) {
+        const char* n = ggml_type_name((ggml_type) t);
+        if (n == nullptr) continue;
+        std::string m(n);
+        for (auto& c : m) c = (char) std::tolower((unsigned char) c);
+        if (m == name) return t;
+    }
+    return -1;
+}
+
+// Random weights (rows of different magnitude) quantized by ggml into the blob layout.
+std::vector<uint8_t> synthetic_blob(const cpu::NativeFmt& f, int seed) {
+    std::vector<uint8_t> blob(f.bytes);
+    std::mt19937 rng(seed);
+    std::normal_distribution<float> nd(0.f, 1.f);
+    auto quant = [&](int type, int64_t rows, int64_t cols, uint8_t* dst) {
+        std::vector<float> w((size_t) (rows * cols));
+        for (int64_t r = 0; r < rows; ++r) {
+            const float scale = 0.02f * (0.5f + (float) (r % 7) / 7.0f);
+            for (int64_t c = 0; c < cols; ++c) w[(size_t) (r * cols + c)] = scale * nd(rng);
+        }
+        const std::vector<float> imatrix((size_t) cols, 1.0f);
+        ggml_quantize_chunk((ggml_type) type, w.data(), dst, 0, rows, cols,
+                            ggml_quantize_requires_imatrix((ggml_type) type) ? imatrix.data() : nullptr);
+    };
+    quant(f.gu_type, FF, H, blob.data());
+    quant(f.gu_type, FF, H, blob.data() + f.up_off);
+    quant(f.d_type, H, FF, blob.data() + f.down_off);
+    return blob;
+}
+
+}  // namespace
+
 int main(int argc, char** argv) {
-    if (argc < 2) { std::fprintf(stderr, "usage: native_expert_parity <shard1.gguf> [layer ...]\n"); return 2; }
-    strata::GgufFile gguf(argv[1]);
-    std::vector<int> layers;
-    for (int i = 2; i < argc; ++i) layers.push_back(std::atoi(argv[i]));
-    if (layers.empty()) layers = {0, 1, 2, 3, 20, 47};
-    const int NT = 3, E = 7;
-    const int64_t H = 2560, FF = 640;
+    if (argc < 2) {
+        std::fprintf(stderr, "usage: native_expert_parity <shard1.gguf> [layer ...]\n"
+                             "       native_expert_parity --synthetic GU/DOWN ...   (ggml type names, e.g. q6_K/q8_0)\n");
+        return 2;
+    }
    int failures = 0;
    cudaStream_t s;
    cudaStreamCreate(&s);
-    for (int l : layers) {
-        const strata::TensorInfo* t[3] = {};
-        const char* roles[3] = {"gate", "up", "down"};
-        for (const auto& ti : gguf.tensors())
-            for (int r = 0; r < 3; ++r)
-                if (ti.name == "blk." + std::to_string(l) + ".ffn_" + roles[r] + "_exps.weight") t[r] = &ti;
-        if (!t[0] || !t[1] || !t[2]) { std::printf("layer %d: no expert tensors\n", l); ++failures; continue; }
-        cpu::NativeFmt f;
-        std::string err;
-        if (!cpu::native_fmt((int) t[0]->type, (int) t[2]->type, H, FF, f, err)) {
-            std::printf("layer %d: %s\n", l, err.c_str()); ++failures; continue;
+    if (std::string(argv[1]) == "--synthetic") {
+        for (int i = 2; i < argc; ++i) {
+            const std::string arg = argv[i];
+            const size_t slash = arg.find('/');
+            const int gu = slash == std::string::npos ? -1 : type_of(arg.substr(0, slash));
+            const int dn = slash == std::string::npos ? -1 : type_of(arg.substr(slash + 1));
+            cpu::NativeFmt f;
+            std::string err;
+            if (gu < 0 || dn < 0 || !cpu::native_fmt(gu, dn, H, FF, f, err) ||
+                !strata::kernels::native_expert_supported(gu, dn, H, FF)) {
+                std::printf("%s: %s\n", arg.c_str(), err.empty() ? "not a supported pair of GGML types" : err.c_str());
+                ++failures;
+                continue;
+            }
+            failures += check_blob(f, synthetic_blob(f, i), i, "synthetic", s);
        }
-        std::vector<uint8_t> blob(f.bytes);
-        std::memcpy(blob.data(), gguf.tensor_data(*t[0]) + (size_t) E * f.up_off, f.up_off);
-        std::memcpy(blob.data() + f.up_off, gguf.tensor_data(*t[1]) + (size_t) E * f.up_off, f.up_off);
-        std::memcpy(blob.data() + f.down_off, gguf.tensor_data(*t[2]) + (size_t) E * (f.bytes - f.down_off),
-                    f.bytes - f.down_off);
-        // (a) the float reference
-        const auto* tg = ggml_get_type_traits((ggml_type) f.gu_type);
-        const auto* td = ggml_get_type_traits((ggml_type) f.d_type);
-        std::vector<float> G((size_t) FF * H), U((size_t) FF * H), D((size_t) H * FF);
-        for (int64_t r = 0; r < FF; ++r) {
-            tg->to_float(blob.data() + r * f.gu_row, G.data() + r * H, H);
-            tg->to_float(blob.data() + f.up_off + r * f.gu_row, U.data() + r * H, H);
+    } else {
+        const strata::GgufModel model(strata::gguf_split_paths(argv[1]));
+        std::vector<int> layers;
+        for (int i = 2; i < argc; ++i) layers.push_back(std::atoi(argv[i]));
+        if (layers.empty()) layers = {0, 1, 2, 3, 20, 47};
+        const int E = 7;
+        for (int l : layers) {
+            const strata::TensorInfo* t[3] = {};
+            const uint8_t* data[3] = {};
+            const char* roles[3] = {"gate", "up", "down"};
+            for (int r = 0; r < 3; ++r) {
+                size_t at = 0;
+                t[r] = model.find("blk." + std::to_string(l) + ".ffn_" + roles[r] + "_exps.weight", &at);
+                if (t[r]) data[r] = model.shard(at).tensor_data(*t[r]);
+            }
+            if (!t[0] || !t[1] || !t[2]) { std::printf("layer %d: no expert tensors\n", l); ++failures; continue; }
+            cpu::NativeFmt f;
+            std::string err;
+            if (!cpu::native_fmt((int) t[0]->type, (int) t[2]->type, H, FF, f, err)) {
+                std::printf("layer %d: %s\n", l, err.c_str()); ++failures; continue;
+            }
+            std::vector<uint8_t> blob(f.bytes);
+            std::memcpy(blob.data(), data[0] + (size_t) E * f.up_off, f.up_off);
+            std::memcpy(blob.data() + f.up_off, data[1] + (size_t) E * f.up_off, f.up_off);
+            std::memcpy(blob.data() + f.down_off, data[2] + (size_t) E * (f.bytes - f.down_off), f.bytes - f.down_off);
+            failures += check_blob(f, blob, l, "layer " + std::to_string(l), s);
        }
-        for (int64_t r = 0; r < H; ++r) td->to_float(blob.data() + f.down_off + r * f.d_row, D.data() + r * FF, FF);
-        std::mt19937 rng(11 + l);
-        std::normal_distribution<float> nd(0.f, 1.f);
-        std::vector<float> x((size_t) NT * H);
-        for (auto& v : x) v = nd(rng);
-        std::vector<float> ref((size_t) NT * H), got_c((size_t) NT * H), got_g((size_t) NT * H);
-        for (int k = 0; k < NT; ++k) {
-            std::vector<float> h(FF);
-            for (int64_t r = 0; r < FF; ++r) {
-                double g = 0, u = 0;
-                for (int64_t i = 0; i < H; ++i) { g += (double) G[r * H + i] * x[k * H + i]; u += (double) U[r * H + i] * x[k * H + i]; }
-                h[r] = (float) (g / (1.0 + std::exp(-g)) * u);
-            }
-            for (int64_t r = 0; r < H; ++r) {
-                double o = 0;
-                for (int64_t i = 0; i < FF; ++i) o += (double) D[r * FF + i] * h[i];
-                ref[k * H + r] = (float) o;
-            }
-        }
-        // (b) the CPU
-        {
-            std::vector<std::vector<uint8_t>> act(NT, std::vector<uint8_t>(cpu::kNativeActBytes));
-            std::vector<std::vector<uint8_t>> hq(NT, std::vector<uint8_t>(cpu::kNativeHBytes));
-            std::vector<std::vector<float>> ff(NT, std::vector<float>(FF));
-            const void* a[NT];
-            float* ffp[NT];
-            const void* hp[NT];
-            float* op[NT];
-            for (int k = 0; k < NT; ++k) {
-                cpu::native_quant_act(f, x.data() + k * H, act[k].data());
-                a[k] = act[k].data();
-                ffp[k] = ff[k].data();
-            }
-            cpu::native_gu_rows(f, blob.data(), a, NT, ffp, 0, (int) FF);
-            if (cpu::iq512_supported(f.gu_type)) {
-                // the AVX-512 rows against ggml's own vec_dot, same Q8_K activations: float-order differences only
-                std::vector<float> g512((size_t) NT * FF), gref((size_t) NT * FF);
-                float* gp[NT];
-                for (int k = 0; k < NT; ++k) gp[k] = g512.data() + k * FF;
-                cpu::iq512_rows(f.gu_type, blob.data(), f.gu_row, (int) H, a, NT, gp, 0, (int) FF);
-                const auto* tc = ggml_get_type_traits_cpu((ggml_type) f.gu_type);
-                for (int k = 0; k < NT; ++k)
-                    for (int64_t r = 0; r < FF; ++r)
-                        tc->vec_dot((int) H, &gref[k * FF + r], 0, blob.data() + r * f.gu_row, 0, a[k], 0, 1);
-                {   // single-thread throughput, the weights cache-resident (compute bound)
-                    float* f1[1] = {ffp[0]};
-                    const int it = 50;
-                    auto t0 = std::chrono::steady_clock::now();
-                    for (int i = 0; i < it; ++i) cpu::iq512_gu_rows(f.gu_type, blob.data(), f.gu_row, f.up_off, (int) H, a, 1, f1, 0, (int) FF);
-                    auto t1 = std::chrono::steady_clock::now();
-                    for (int i = 0; i < it; ++i)
-                        for (int64_t r = 0; r < FF; ++r) {
-                            float gg, uu;
-                            tc->vec_dot((int) H, &gg, 0, blob.data() + r * f.gu_row, 0, a[0], 0, 1);
-                            tc->vec_dot((int) H, &uu, 0, blob.data() + f.up_off + r * f.gu_row, 0, a[0], 0, 1);
-                        }
-                    auto t2 = std::chrono::steady_clock::now();
-                    const double us512 = std::chrono::duration<double, std::micro>(t1 - t0).count() / it;
-                    const double usg = std::chrono::duration<double, std::micro>(t2 - t1).count() / it;
-                    std::printf("          gate+up one token one thread: avx512 %.0f us (%.2f GB/s), ggml %.0f us (%.2f GB/s)\n",
-                                us512, 2.0 * f.up_off / us512 / 1e3, usg, 2.0 * f.up_off / usg / 1e3);
-                    auto t3 = std::chrono::steady_clock::now();
-                    for (int i = 0; i < it; ++i) cpu::iq512_gu_rows(f.gu_type, blob.data(), f.gu_row, f.up_off, (int) H, a, NT, ffp, 0, (int) FF);
-                    auto t4 = std::chrono::steady_clock::now();
-                    std::printf("          gate+up %d tokens one thread: avx512 %.0f us vs ggml %d x %.0f us\n", NT,
-                                std::chrono::duration<double, std::micro>(t4 - t3).count() / it, NT, usg);
-                }
-                std::printf("          %s AVX-512 gate rows vs ggml vec_dot: rel %.2e\n",
-                            ggml_type_name((ggml_type) f.gu_type), rel(g512, gref));
-            }
-            for (int k = 0; k < NT; ++k) {
-                cpu::native_quant_h(f, ff[k].data(), hq[k].data());
-                hp[k] = hq[k].data();
-                op[k] = got_c.data() + k * H;
-            }
-            cpu::native_down_rows(f, blob.data(), hp, NT, op, 0, (int) H);
-            if (f.d_type == 42) {
-                // (b2) the AVX-512 GGUF-layout Q2_0 kernel the pool uses for Q2_0 down projections
-                std::vector<cpu::ActQ> a2(NT);
-                const cpu::ActQ* ap[NT];
-                std::vector<float> alt((size_t) NT * H);
-                float* altp[NT];
-                for (int k = 0; k < NT; ++k) {
-                    cpu::act_quant_q8_1(ff[k].data(), (int) FF, a2[k]);
-                    ap[k] = &a2[k];
-                    altp[k] = alt.data() + k * H;
-                }
-                cpu::q2_0_gguf_rows_multi(blob.data() + f.down_off, f.d_row, (int) (FF / 64), ap, NT, altp, 0, (int) H);
-                std::printf("          q2_0 AVX-512 down vs ggml down: rel %.2e\n", rel(alt, got_c));
-            }
-        }
-        // (c) the GPU: one group holding the NT entries
-        {
-            const auto L = strata::kernels::native_expert_layout(f.gu_type, f.d_type, H, FF);
-            void *dblob, *dx, *dxq, *dscr;
-            float* dout;
-            unsigned long long* dptr;
-            int32_t *dstart, *dn, *ddst, *dtok;
-            cudaMalloc(&dblob, blob.size());
-            cudaMalloc(&dx, x.size() * 4);
-            cudaMalloc(&dxq, (size_t) NT * H / 32 * 36);
-            cudaMalloc(&dscr, strata::kernels::native_expert_scratch_bytes(NT, FF));
-            cudaMalloc((void**) &dout, (size_t) NT * H * 4);
-            cudaMalloc((void**) &dptr, 8);
-            cudaMalloc((void**) &dstart, 8);
-            cudaMalloc((void**) &dn, 4);
-            cudaMalloc((void**) &ddst, NT * 4);
-            cudaMalloc((void**) &dtok, NT * 4);
-            cudaMemcpy(dblob, blob.data(), blob.size(), cudaMemcpyHostToDevice);
-            cudaMemcpy(dx, x.data(), x.size() * 4, cudaMemcpyHostToDevice);
-            const unsigned long long p = (unsigned long long) dblob;
-            const int32_t st[2] = {0, NT}, one = 1, idx[NT] = {0, 1, 2};
-            cudaMemcpy(dptr, &p, 8, cudaMemcpyHostToDevice);
-            cudaMemcpy(dstart, st, 8, cudaMemcpyHostToDevice);
-            cudaMemcpy(dn, &one, 4, cudaMemcpyHostToDevice);
-            cudaMemcpy(ddst, idx, NT * 4, cudaMemcpyHostToDevice);
-            cudaMemcpy(dtok, idx, NT * 4, cudaMemcpyHostToDevice);
-            strata::kernels::quantize_q8_1_rows((const float*) dx, NT, H, dxq, s);
-            strata::kernels::native_expert_grouped(L, dptr, dstart, dn, ddst, dtok, 1, NT, dxq, dscr, dout, s);
-            cudaStreamSynchronize(s);
-            cudaMemcpy(got_g.data(), dout, got_g.size() * 4, cudaMemcpyDeviceToHost);
-            cudaFree(dblob); cudaFree(dx); cudaFree(dxq); cudaFree(dscr); cudaFree(dout); cudaFree(dptr);
-            cudaFree(dstart); cudaFree(dn); cudaFree(ddst); cudaFree(dtok);
-        }
-        const double ec = rel(got_c, ref), eg = rel(got_g, ref), ecg = rel(got_c, got_g);
-        const bool ok = ec < 3e-2 && eg < 3e-2 && std::isfinite(ec) && std::isfinite(eg);
-        std::printf("layer %2d  %-8s/%-7s blob %8zu  cpu rel %.2e  gpu rel %.2e  cpu-gpu %.2e  %s\n", l,
-                    ggml_type_name((ggml_type) f.gu_type), ggml_type_name((ggml_type) f.d_type), f.bytes, ec, eg, ecg,
-                    ok ? "ok" : "FAIL");
-        if (!ok) ++failures;
    }
    std::printf("native_expert_parity: %d failures\n", failures);
    return failures ? 1 : 0;
@@ -1,4 +1,4 @@
-// src/kernels/ngram.cpp - P2.S4: the PLE n-gram hash and the IQ4_NL table read.
+// src/kernels/ngram.cpp - P2.S4: the PLE n-gram hash and the table read (IQ4_NL or Q8_0).
 //
 // See include/strata/kernels/ngram.hpp for the semantics, the rival readings, and the note on MADV_RANDOM.
 #include "strata/kernels/ngram.hpp"
@@ -107,11 +107,37 @@ void iq4nl_dequant_row(const uint8_t* row, float* out160) {
    }
 }

+void ple_dequant_row(int type, const uint8_t* row, float* out160) {
+    if (type != 8) {
+        iq4nl_dequant_row(row, out160);
+        return;
+    }
+    for (int b = 0; b < PLE_HEAD_DIM / 32; ++b) {                     // Q8_0: fp16 scale + 32 int8
+        const uint8_t* blk = row + (size_t) b * 34;
+        uint16_t dbits;
+        std::memcpy(&dbits, blk, 2);
+        const float d = f32_from_f16(dbits);
+        for (int j = 0; j < 32; ++j) out160[b * 32 + j] = d * (float) (int8_t) blk[2 + j];
+    }
+}
+
+std::string ple_table_shard(const std::vector<std::string>& shards) {
+    for (const auto& path : shards) {
+        try {
+            if (GgufFile(path).find("per_layer_token_embd.weight") != nullptr) return path;
+        } catch (const std::exception&) {
+        }
+    }
+    return {};
+}
+
 // ---------------------------------------------------------------------------------------------------
 struct PleTable::Impl {
    GgufFile* file = nullptr;
    const uint8_t* data = nullptr;
    uint64_t n_rows = 0;
+    int type = 20;                        // IQ4_NL or Q8_0
+    uint32_t row_bytes = PLE_ROW_BYTES;
    mutable uint64_t bytes_read = 0;
    // Direct mode (plan v0.3 P2): the mapping above is released after the header parse and every row comes
    // from an unbuffered SSD read into `raw`.
@@ -120,7 +146,7 @@ struct PleTable::Impl {
    strata::ngram::PleReader::Ticket ticket;
    bool pending = false;
    uint32_t rows[PLE_N_HEADS] = {};
-    uint8_t raw[PLE_N_HEADS * PLE_ROW_BYTES] = {};
+    uint8_t raw[PLE_N_HEADS * PLE_ROW_BYTES_Q8_0] = {};
 };

 PleTable::PleTable() : impl_(new Impl) {}
@@ -150,11 +176,13 @@ bool PleTable::open(const std::string& gguf_path, std::string& err, const PleIoO
        close();
        return false;
    }
-    if (std::strcmp(t->type_name(), "IQ4_NL") != 0) {
-        err = std::string("per_layer_token_embd.weight is ") + t->type_name() + ", not IQ4_NL";
+    if (t->type != 20 && t->type != 8) {
+        err = std::string("per_layer_token_embd.weight is ") + t->type_name() + ", not IQ4_NL or Q8_0";
        close();
        return false;
    }
+    impl_->type = (int) t->type;
+    impl_->row_bytes = t->type == 8 ? PLE_ROW_BYTES_Q8_0 : PLE_ROW_BYTES;
    impl_->n_rows = t->shape[1];
    impl_->data = impl_->file->tensor_data(*t);

@@ -162,16 +190,21 @@ bool PleTable::open(const std::string& gguf_path, std::string& err, const PleIoO
    // the offset within the GGUF's DATA SECTION: the file's first 192 bytes are a header, and reading at 0
    // would decode the header plus 192 bytes of shifted rows - still plausible IQ4_NL, and wrong for every
    // row.  `GgufFile` parses the header, so `tensor_data` is already correct; this asserts the tensor
-    // exactly fills the file from there, which is what makes the whole arrangement checkable rather than
-    // assumed.  A wrong data offset would leave a different remainder.
-    const uint64_t need = impl_->n_rows * (uint64_t) PLE_ROW_BYTES;
-    const uint64_t have = impl_->file->file_size() - impl_->file->data_start();
-    if (need != have) {
+    // exactly fills its span - up to the next tensor or the end of the file, less the alignment padding -
+    // which is what makes the whole arrangement checkable rather than assumed.  A wrong data offset would
+    // leave a different remainder.  (The table fills the ISTA shard 2 alone; the unsloth shard 2 holds other
+    // tensors too.)
+    const uint64_t need = impl_->n_rows * (uint64_t) impl_->row_bytes;
+    uint64_t end = impl_->file->file_size() - impl_->file->data_start();
+    for (const auto& o : impl_->file->tensors())
+        if (o.offset > t->offset && o.offset < end) end = o.offset;
+    const uint64_t have = end - t->offset;
+    if (need > have || have - need >= impl_->file->alignment()) {
        char buf[256];
        std::snprintf(buf, sizeof buf,
-                      "PLE table size mismatch: %llu rows x %d B = %llu, but the file holds %llu from "
-                      "data_start %llu",
-                      (unsigned long long) impl_->n_rows, PLE_ROW_BYTES, (unsigned long long) need,
+                      "PLE table size mismatch: %llu rows x %d B = %llu, but its span in the file is %llu B "
+                      "(data_start %llu)",
+                      (unsigned long long) impl_->n_rows, (int) impl_->row_bytes, (unsigned long long) need,
                      (unsigned long long) have, (unsigned long long) impl_->file->data_start());
        err = buf;
        close();
@@ -185,7 +218,8 @@ bool PleTable::open(const std::string& gguf_path, std::string& err, const PleIoO
        delete impl_->file;
        impl_->file = nullptr;
        impl_->data = nullptr;
-        if (!impl_->reader.open(gguf_path, table_offset, n_rows, io.max_inflight, io.cache_rows, err, io.io_thread)) {
+        if (!impl_->reader.open(gguf_path, table_offset, n_rows, io.max_inflight, io.cache_rows, err, io.io_thread,
+                                impl_->row_bytes)) {
            close();
            return false;
        }
@@ -212,23 +246,23 @@ uint64_t PleTable::bytes_read() const { return impl_->bytes_read; }

 void PleTable::read_row(uint32_t row, float* out160) const {
    if (impl_->mode == PleIo::Direct && impl_->reader.is_open()) {
-        uint8_t raw[PLE_ROW_BYTES];
+        uint8_t raw[PLE_ROW_BYTES_Q8_0];
        std::string err;
        const auto t = impl_->reader.issue(&row, 1, raw);
        if (!impl_->reader.collect(t, err)) {
            std::memset(out160, 0, (size_t) PLE_HEAD_DIM * sizeof(float));
            return;
        }
-        iq4nl_dequant_row(raw, out160);
-        impl_->bytes_read += PLE_ROW_BYTES;
+        ple_dequant_row(impl_->type, raw, out160);
+        impl_->bytes_read += impl_->row_bytes;
        return;
    }
    if (impl_->data == nullptr || row >= impl_->n_rows) {
        std::memset(out160, 0, (size_t) PLE_HEAD_DIM * sizeof(float));
        return;
    }
-    iq4nl_dequant_row(impl_->data + (size_t) row * PLE_ROW_BYTES, out160);
-    impl_->bytes_read += PLE_ROW_BYTES;
+    ple_dequant_row(impl_->type, impl_->data + (size_t) row * impl_->row_bytes, out160);
+    impl_->bytes_read += impl_->row_bytes;
 }

 bool PleTable::issue(const uint32_t* rows16) {
@@ -245,8 +279,8 @@ bool PleTable::issue(const uint32_t* rows16) {
        ULONG_PTR n = 0;
        for (int h = 0; h < PLE_N_HEADS; ++h) {
            if (rows16[h] >= impl_->n_rows) continue;
-            ranges[n].VirtualAddress = (PVOID) (impl_->data + (size_t) rows16[h] * PLE_ROW_BYTES);
-            ranges[n].NumberOfBytes = PLE_ROW_BYTES;
+            ranges[n].VirtualAddress = (PVOID) (impl_->data + (size_t) rows16[h] * impl_->row_bytes);
+            ranges[n].NumberOfBytes = impl_->row_bytes;
            ++n;
        }
        if (n > 0) (void) PrefetchVirtualMemory(GetCurrentProcess(), n, ranges, 0);
@@ -262,8 +296,8 @@ bool PleTable::collect(float* out2560, std::string& err) {
    if (impl_->mode == PleIo::Direct) {
        if (!impl_->reader.collect(impl_->ticket, err)) return false;
        for (int h = 0; h < PLE_N_HEADS; ++h)
-            iq4nl_dequant_row(impl_->raw + (size_t) h * PLE_ROW_BYTES, out2560 + (size_t) h * PLE_HEAD_DIM);
-        impl_->bytes_read += (uint64_t) PLE_N_HEADS * PLE_ROW_BYTES;
+            ple_dequant_row(impl_->type, impl_->raw + (size_t) h * impl_->row_bytes, out2560 + (size_t) h * PLE_HEAD_DIM);
+        impl_->bytes_read += (uint64_t) PLE_N_HEADS * impl_->row_bytes;
        return true;
    }
    for (int h = 0; h < PLE_N_HEADS; ++h) read_row(impl_->rows[h], out2560 + (size_t) h * PLE_HEAD_DIM);
@@ -274,11 +308,12 @@ bool PleTable::gather_batch(const uint32_t* rows, size_t n_tokens, float* out, s
    if (impl_->pending) { err = "PleTable::gather_batch while a token is in flight"; return false; }
    const size_t n = n_tokens * (size_t) PLE_N_HEADS;
    if (impl_->mode == PleIo::Direct) {
-        std::vector<uint8_t> raw(n * PLE_ROW_BYTES);
+        std::vector<uint8_t> raw(n * impl_->row_bytes);
        const auto ticket = impl_->reader.issue(rows, n, raw.data());
        if (!impl_->reader.collect(ticket, err)) return false;
-        for (size_t i = 0; i < n; ++i) iq4nl_dequant_row(raw.data() + i * PLE_ROW_BYTES, out + i * PLE_HEAD_DIM);
-        impl_->bytes_read += (uint64_t) n * PLE_ROW_BYTES;
+        for (size_t i = 0; i < n; ++i)
+            ple_dequant_row(impl_->type, raw.data() + i * impl_->row_bytes, out + i * PLE_HEAD_DIM);
+        impl_->bytes_read += (uint64_t) n * impl_->row_bytes;
        return true;
    }
    for (size_t i = 0; i < n; ++i) read_row(rows[i], out + i * PLE_HEAD_DIM);
@@ -332,8 +367,8 @@ void PleTable::gather(const uint32_t* rows16, float* out2560) const {
        for (int h = 0; h < PLE_N_HEADS; ++h) {
            // Out-of-range rows are handled by `read_row` as zeros and have no address to prefetch.
            if (rows16[h] >= impl_->n_rows) continue;
-            ranges[n].VirtualAddress = (PVOID) (impl_->data + (size_t) rows16[h] * PLE_ROW_BYTES);
-            ranges[n].NumberOfBytes = PLE_ROW_BYTES;
+            ranges[n].VirtualAddress = (PVOID) (impl_->data + (size_t) rows16[h] * impl_->row_bytes);
+            ranges[n].NumberOfBytes = impl_->row_bytes;
            ++n;
        }
        if (n > 0) (void) PrefetchVirtualMemory(GetCurrentProcess(), n, ranges, 0);
@@ -33,14 +33,16 @@ constexpr uint32_t EMPTY = 0xFFFFFFFFu;
 struct RowCache {
    uint64_t sets = 0;
    std::vector<uint32_t> keys;     // sets * WAYS
-    std::vector<uint8_t> data;      // sets * WAYS * ROW_BYTES
+    std::vector<uint8_t> data;      // sets * WAYS * row_bytes
    std::vector<uint8_t> next;      // per-set replacement pointer
    uint64_t used = 0;
+    uint32_t row_bytes = ROW_BYTES;

-    void init(uint64_t rows) {
+    void init(uint64_t rows, uint32_t bytes) {
        sets = rows / WAYS;
+        row_bytes = bytes;
        keys.assign(sets * WAYS, EMPTY);
-        data.assign(sets * WAYS * ROW_BYTES, 0);
+        data.assign(sets * WAYS * row_bytes, 0);
        next.assign(sets, 0);
        used = 0;
    }
@@ -52,7 +54,7 @@ struct RowCache {
        if (sets == 0) return nullptr;
        const uint64_t s = mix(row) % sets;
        for (uint32_t w = 0; w < WAYS; ++w)
-            if (keys[s * WAYS + w] == row) return &data[(s * WAYS + w) * ROW_BYTES];
+            if (keys[s * WAYS + w] == row) return &data[(s * WAYS + w) * row_bytes];
        return nullptr;
    }
    void insert(uint32_t row, const uint8_t* bytes) {
@@ -62,7 +64,7 @@ struct RowCache {
        next[s] = (uint8_t) ((w + 1) % WAYS);
        if (keys[s * WAYS + w] == EMPTY) ++used;
        keys[s * WAYS + w] = row;
-        std::memcpy(&data[(s * WAYS + w) * ROW_BYTES], bytes, ROW_BYTES);
+        std::memcpy(&data[(s * WAYS + w) * row_bytes], bytes, row_bytes);
    }
 };

@@ -94,6 +96,7 @@ struct PleReader::Impl {
    DirectFile file;
    uint64_t table_offset = 0;
    uint64_t n_rows = 0;
+    uint32_t row_bytes = ROW_BYTES;
    uint32_t max_inflight = 0;
    uint8_t* slab = nullptr;              // max_inflight slots of 2 pages
    std::vector<uint32_t> free_slots;
@@ -150,11 +153,11 @@ struct PleReader::Impl {
        stats.bytes += c.bytes;
        const uint8_t* buf = slot_buf(s);
        for (const Use& u : j.uses) {
-            if (u.in_page + ROW_BYTES > c.bytes) {
+            if (u.in_page + row_bytes > c.bytes) {
                error = "PleReader: short read inside the table";
                return false;
            }
-            std::memcpy(u.dst, buf + u.in_page, ROW_BYTES);
+            std::memcpy(u.dst, buf + u.in_page, row_bytes);
            cache.insert(u.row, buf + u.in_page);
        }
        auto it = tickets.find(j.ticket);
@@ -234,24 +237,26 @@ PleReader::~PleReader() {
 }

 bool PleReader::open(const std::string& path, uint64_t table_offset, uint64_t n_rows, uint32_t max_inflight,
-                     uint64_t cache_rows, std::string& err, bool io_thread) {
+                     uint64_t cache_rows, std::string& err, bool io_thread, uint32_t row_bytes) {
    close();
    if (max_inflight == 0 || max_inflight > 1024) { err = "PleReader: max_inflight must be 1..1024"; return false; }
+    if (row_bytes == 0 || row_bytes > PAGE) { err = "PleReader: a row must fit a page"; return false; }
    if (!impl_->file.open(path, err)) return false;
-    if (table_offset + n_rows * (uint64_t) ROW_BYTES > impl_->file.size()) {
+    if (table_offset + n_rows * (uint64_t) row_bytes > impl_->file.size()) {
        err = "PleReader: the table extends past the end of " + path;
        close();
        return false;
    }
    impl_->table_offset = table_offset;
    impl_->n_rows = n_rows;
+    impl_->row_bytes = row_bytes;
    impl_->max_inflight = max_inflight;
    impl_->slab = (uint8_t*) DirectFile::alloc_aligned((size_t) max_inflight * 2 * PAGE);
    if (impl_->slab == nullptr) { err = "PleReader: cannot allocate read buffers"; close(); return false; }
    impl_->inflight.assign(max_inflight, Job{});
    impl_->free_slots.clear();
    for (uint32_t s = max_inflight; s-- > 0;) impl_->free_slots.push_back(s);
-    impl_->cache.init(cache_rows);
+    impl_->cache.init(cache_rows, row_bytes);
    impl_->error.clear();
    reset_stats();
    impl_->stop = false;
@@ -291,7 +296,7 @@ void PleReader::close() {
    m.delayed.clear();
    m.inflight.clear();
    m.free_slots.clear();
-    m.cache.init(0);
+    m.cache.init(0, ROW_BYTES);
    m.threaded = false;
    m.stop = false;
 }
@@ -308,20 +313,20 @@ PleReader::Ticket PleReader::issue(const uint32_t* rows, size_t n, uint8_t* out_
    std::unordered_map<uint64_t, size_t> by_page;     // aligned offset -> index in `jobs`
    std::vector<Job> jobs;
    for (size_t i = 0; i < n; ++i) {
-        uint8_t* dst = out_raw + i * ROW_BYTES;
+        uint8_t* dst = out_raw + i * m.row_bytes;
        ++m.stats.requests;
        if (rows[i] >= m.n_rows) {
-            std::memset(dst, 0, ROW_BYTES);
+            std::memset(dst, 0, m.row_bytes);
            continue;
        }
        if (const uint8_t* hit = m.cache.find(rows[i])) {
-            std::memcpy(dst, hit, ROW_BYTES);
+            std::memcpy(dst, hit, m.row_bytes);
            ++m.stats.cache_hits;
            continue;
        }
-        const uint64_t at = m.table_offset + (uint64_t) rows[i] * ROW_BYTES;
+        const uint64_t at = m.table_offset + (uint64_t) rows[i] * m.row_bytes;
        const uint64_t first = at / PAGE * PAGE;
-        const uint32_t length = (uint32_t) ((at + ROW_BYTES - 1) / PAGE * PAGE - first + PAGE);
+        const uint32_t length = (uint32_t) ((at + m.row_bytes - 1) / PAGE * PAGE - first + PAGE);
        auto f = by_page.find(first);
        if (f != by_page.end()) {
            Job& j = jobs[f->second];
@@ -14,6 +14,7 @@
 // AND IT IS PHASE 2, so hit rate is `h = 0` and the number it prints is slow on purpose
 // (`phase-2-correct-engine.md:5-9`).  What it is FOR is the honest tok/s figure and the logit dump.

+#include "strata/artifact/gguf_split.hpp"
 #include "strata/core/expert_cache.hpp"
 #include "strata/core/expert_source.hpp"
 #include "strata/core/layer.hpp"
@@ -27,7 +28,9 @@
 #include "strata/kernels/s2_expert_grouped.hpp"
 #include "strata/kernels/sampler.hpp"
 #include "strata/kernels/shared_expert.hpp"
+#include "strata/kernels/iq_kernels.hpp"
 #include "strata/kernels/native_moe.hpp"
+#include "strata/kernels/native_mmvq.hpp"
 #include "strata/kernels/native_gdn.hpp"
 #include "strata/kernels/native_router.hpp"
 #include "strata/kernels/native_qsa.hpp"
@@ -251,7 +254,8 @@ void usage() {
                 "  --pack DIR           the pack directory (default pack/full)\n"
                 "  --tokens LIST        the prompt as comma-separated token IDS (required)\n"
                 "  --tokens-file PATH   pretokenized prompt, commas or whitespace (alternative to --tokens)\n"
-                 "  --ple-gguf PATH      required PLE table (original second GGUF shard)\n"
+                 "  --ple-gguf PATH      the PLE table: the model shard holding per_layer_token_embd (with --native,\n"
+                 "                       the --native model's shard that holds it by default)\n"
                 "  --no-ple             explicit diagnostic ablation of the PLE layer\n"
                 "  --ple-io direct|mmap  n-gram table reads (plan v0.3 P2). direct (default): unbuffered SSD\n"
                 "                       reads, the table never enters RAM or the file cache; mmap: A/B arm\n"
@@ -277,11 +281,12 @@ void usage() {
                 "  --native-ple-postops  experimental pinned PLE postprojection arithmetic\n"
                 "  --native-router       experimental pinned CUDA 512-expert top-10 routing\n"
                 "  --cpu-oracle-q8-0     experimental pinned CPU expert quantization and dot reduction\n"
-                 "  --native SHARD1      every full-context native path at once (plan v0.3 P1): stream-token,\n"
+                 "  --native SHARD1      the model's first GGUF shard (the others of a split model are found by\n"
+                 "                       name) and every full-context native path at once (plan v0.3 P1): stream-token,\n"
                 "                       GR MMVF, BF16, head, dense + PLE key, MoE combine, GDN, router, QSA,\n"
                 "                       indexer, RoPE, PLE postops, and the CPU q8_0 contract unless the\n"
                 "                       expert cache is on. Individual --native-* flags stay for A/B.\n"
-                 "  --native-head-gguf PATH  native Q5_K head from model shard 1; requires --stream-token\n"
+                 "  --native-head-gguf PATH  native head from the model with this first shard; requires --stream-token\n"
                 "  --native-dense-gguf PATH native GDN/QSA/shared projections; repeat for each source model shard\n"
                 "  --expert-cache-cpu-order  experimental GPU expert reduction matching CPU order\n"
                 "  --max-new N          tokens to generate (default 16)\n"
@@ -312,8 +317,6 @@ void usage() {
                 "  --prefill CHUNK      batched prompt processing in chunks of CHUNK tokens (needs --native)\n"
                 "  --no-pool            skip the CPU expert pool (the GPU-only floor)\n"
                 "  --sync-every-layer   debug: synchronise after every layer\n"
-                 "  --ple-gguf PATH      the n-gram/PLE shard.  WITHOUT IT LAYER 1's PLE IS SILENTLY SKIPPED,\n"
-                 "                       which changes every number downstream - pass it for any real run\n"
                 "  --dump-mixed PATH    write the post-attention residual (n_embd, f32)\n"
                 "  --stage-timing       per-stage KERNEL-COUNT shares.  NOT a time profile: an uncaptured\n"
                 "                       event interval includes host gaps, so run with --gpu-only-full first\n"
@@ -604,9 +607,22 @@ int main(int argc, char** argv) {
    }
    strata::core::qsa_set_kv_int8(o.kv == "int8");
    strata::core::layer_set_shared_early(!o.shared_late);
+    std::vector<std::string> native_shards;   // the --native model's GGUF shards
    if (!o.native_preset.empty()) {
-        if (o.no_ple || o.ple_gguf.empty()) {
-            std::fprintf(stderr, "strata generate: --native requires --ple-gguf (the PLE key is native too)\n");
+        if (o.no_ple) {
+            std::fprintf(stderr, "strata generate: --native requires the PLE (the PLE key is native too)\n");
+            return 2;
+        }
+        try {
+            native_shards = strata::gguf_split_paths(o.native_preset);
+        } catch (const std::exception& e) {
+            std::fprintf(stderr, "strata generate: --native: %s\n", e.what());
+            return 2;
+        }
+        if (o.ple_gguf.empty()) o.ple_gguf = strata::kernels::ple_table_shard(native_shards);
+        if (o.ple_gguf.empty()) {
+            std::fprintf(stderr, "strata generate: no shard of the --native model holds per_layer_token_embd; "
+                                 "name the file with --ple-gguf\n");
            return 2;
        }
        o.stream_token = true;
@@ -615,7 +631,11 @@ int main(int argc, char** argv) {
        o.native_ple_key = o.native_moe_combine = o.native_gdn = o.native_router = true;
        o.native_qsa = o.native_qsa_indexer = o.native_rope = o.native_ple_postops = true;
        if (o.native_head_gguf.empty()) o.native_head_gguf = o.native_preset;
-        if (o.native_dense_gguf.empty()) o.native_dense_gguf = {o.native_preset, o.ple_gguf};
+        if (o.native_dense_gguf.empty()) {
+            o.native_dense_gguf = native_shards;
+            if (std::find(native_shards.begin(), native_shards.end(), o.ple_gguf) == native_shards.end())
+                o.native_dense_gguf.push_back(o.ple_gguf);
+        }
        // Plan v0.3 (24 Sep): the CPU experts stay on the VNNI kernel.  The llama.cpp-CPU-exact q8_0 contract
        // cost 27.0 vs 17.2 ms/token of pool time and G-C does not need it; `--cpu-oracle-q8-0` still selects it.
    }
@@ -689,6 +709,15 @@ int main(int argc, char** argv) {
        }
    }
    const bool native_pack = strata::kernels::cpu::expert_layout().native;
+    if (native_pack)
+        for (int64_t l = 0; l < strata::kernels::cpu::expert_layout().n_layers; ++l) {
+            const auto& f = strata::kernels::cpu::expert_layout().fmt[(size_t) l];
+            if (!strata::kernels::native_expert_supported(f.gu_type, f.d_type, f.n_embd, f.n_ff)) {
+                std::fprintf(stderr, "strata generate: layer %lld's experts (GGML types %d / %d) have no GPU kernels\n",
+                             (long long) l, f.gu_type, f.d_type);
+                return 1;
+            }
+        }
    // plan v0.3 P6: the PCIe share of the missed experts, measured per kind of pack (the paper, finding on PCIe)
    if (o.pcie_frac < 0.0) o.pcie_frac = native_pack ? 0.55 : 0.2;
    // the canonical Q2_0 pack's CPU kernels are AVX-512 only; a native pack runs on AVX2 CPUs as well
@@ -704,7 +733,7 @@ int main(int argc, char** argv) {
            return 2;
        }
        const strata::core::ModelGeometry g0;
-        if (!native_embed.load(o.native_preset, g0.n_embd, 248320, err)) {
+        if (!native_embed.load(native_shards, g0.n_embd, 248320, err)) {
            std::fprintf(stderr, "strata generate: %s\n", err.c_str());
            return 1;
        }
@@ -900,7 +929,8 @@ int main(int argc, char** argv) {
            ss.ple.w.key_scales = (const float*) ((const uint8_t*) wk->data + wk->codes_bytes);
        }
        if (o.native_ple_key && wk->quantized()) {
-            if (!wk->native_data || wk->native_type != 42 || !wk->native_q8_1) {
+            // Q2_0 in the Q2_0 file, Q8_0 in UD-Q4_K_XL
+            if (!wk->native_data || !strata::kernels::native_mmvq_supported(wk->native_type) || !wk->native_q8_1) {
                std::fprintf(stderr, "strata generate: native PLE key is absent or incompatible\n");
                return 1;
            }
@@ -977,7 +1007,7 @@ int main(int argc, char** argv) {
        std::fprintf(stderr, "strata generate: experts via mmap (--mmap-experts; the A/B arm of R2.1)\n");
        srcp = &src;
    } else {
-        arena_src.set_gguf(o.native_preset);   // plan v0.3 P6: a native pack may take its experts from shard 1
+        arena_src.set_gguf(native_shards);   // plan v0.3 P6: a native pack may take its experts from the GGUF
        if (!arena_src.open(o.pack, g.n_layers, g.n_expert, /*threads=*/6, err)) {
            std::fprintf(stderr, "strata generate: %s\n", err.c_str());
            return 1;
@@ -1295,11 +1325,18 @@ int main(int argc, char** argv) {
    const int64_t n_vocab = wo->ne1;
    strata::core::NativeHead native_head;
    if (!o.native_head_gguf.empty()) {
-        if (!native_head.load(o.native_head_gguf, g.n_embd, n_vocab, err)) {
+        std::vector<std::string> head_shards;
+        try {
+            head_shards = strata::gguf_split_paths(o.native_head_gguf);
+        } catch (const std::exception& e) {
+            std::fprintf(stderr, "strata generate: native head: %s\n", e.what());
+            return 1;
+        }
+        if (!native_head.load(head_shards, g.n_embd, n_vocab, err)) {
            std::fprintf(stderr, "strata generate: %s\n", err.c_str());
            return 1;
        }
-        std::fprintf(stderr, "strata generate: experimental native Q5_K head, %llu bytes\n",
+        std::fprintf(stderr, "strata generate: native head, type %d, %llu bytes\n", native_head.type(),
                     (unsigned long long) native_head.weight_bytes());
    }
    std::vector<float> logits((size_t) n_vocab);
@@ -1,19 +1,22 @@
-"""tools/iq_pack.py - plan v0.3 P6: a native pack for any of the model files (Q2_0, IQ2_XS, IQ3_XXS).
+"""tools/iq_pack.py - plan v0.3 P6: a native pack for a model file whose tensors the engine reads in GGUF form
+(Q2_0, IQ2_XS, IQ3_XXS, UD-Q4_K_XL).

-    python tools/iq_pack.py --gguf <model>-00001-of-00002.gguf --out pack/iq3_xxs            (standalone)
+    python tools/iq_pack.py --gguf <model>-00001-of-0000N.gguf --out pack/iq3_xxs            (standalone)
    python tools/iq_pack.py --gguf <model>-00001-of-00002.gguf --base pack/full --out ...    (share dense.bin)

-The i-quant experts cannot be re-expressed in the Q2_0 pack form, so this pack keeps every quantized tensor in
-its GGUF form:
+The other shards of a split model are found by name.  The experts cannot be re-expressed in the Q2_0 pack form,
+so this pack keeps every quantized tensor in its GGUF form:

-  experts.bin          per layer, 512 blobs of [gate rows | up rows | down rows], the raw GGUF slices.  Blob
-                       size is per layer (the files mix IQ1_M ... IQ3_S gate/up and Q2_0 / IQ4_NL down).
-  native_experts.txt   one line per layer: layer gu_type d_type offset blob_bytes
+  native_experts.txt   one line per layer: layer gu_type d_type offset blob_bytes.  The engine fills its expert
+                       arena from the GGUF (tensors found by name); `offset` is the layer's place in that arena.
+  experts.bin          optional (--experts-bin): the arena as a file, per layer 512 blobs of [gate rows | up rows |
+                       down rows], the raw GGUF slices.  Blob size is per layer (the files mix formats).
  index.txt            the table the engine loads.  Quantized dense tensors, token_embd and output are served
                       natively from the GGUF by the engine (--native): their rows carry shape only.
-  dense.bin            standalone: the BF16/F16/F32 tensors exactly as the GGUF stores them (index kinds 4/5/2).
+  dense.bin            standalone: every other tensor in the float form the engine takes (FORM below; index kinds
+                       4/5/2 = BF16/F16/F32), converted when the file stores it otherwise.
                       With --base: the base (Q2_0) pack's dense.bin, hard-linked - the float tensors are
-                       byte-identical in all three model files (checked) - plus extra.bin for tensors that are
+                       byte-identical in all three ISTA model files (checked) - plus extra.bin for tensors that are
                       float here but quantized in the base pack (blk.1.ple_key).
  tokenizer/           exported from the GGUF (tools/strata_tokenizer.py).
 """
@@ -23,6 +26,7 @@ import argparse
 import json
 import os
 import pathlib
+import re
 import shutil
 import subprocess
 import sys
@@ -37,6 +41,97 @@ FLOAT = {"BF16", "F32", "F16"}
 ROLES = ("gate", "up", "down")
 N_EXPERT = 512
 ALIGN = 64
+KIND = {"BF16": "4", "F16": "5", "F32": "2"}
+
+# The float form the engine takes for each tensor it does not read from the GGUF (name without `blk.N.`).  A file
+# that stores one otherwise is converted: UD-Q4_K_XL keeps the router, SSM gates and injections in F32 (their
+# values are BF16's, so the conversion is exact), the hyper-connection projections and the PLE value in Q8_0
+# (re-rounded to BF16) and the PLE conv in F32.
+FORM = {
+    "ffn_gate_inp.weight": "BF16", "ffn_gate_inp_shexp.weight": "BF16",
+    "hc_attn_down.weight": "BF16", "hc_attn_up.weight": "BF16", "hc_attn_inject.weight": "BF16",
+    "hc_ffn_down.weight": "BF16", "hc_ffn_up.weight": "BF16", "hc_ffn_inject.weight": "BF16",
+    "output_hc_down.weight": "BF16", "output_hc_up.weight": "BF16",
+    "indexer.k_proj.weight": "BF16", "indexer.q_proj.weight": "BF16",
+    "ssm_alpha.weight": "BF16", "ssm_beta.weight": "BF16",
+    "ple_key.weight": "BF16", "ple_value.weight": "BF16", "ple_conv1d.weight": "F16",
+    "attn_q_norm.weight": "F32", "attn_k_norm.weight": "F32", "hc_attn_norm.weight": "F32",
+    "hc_ffn_norm.weight": "F32", "output_hc_norm.weight": "F32", "indexer.q_norm.weight": "F32",
+    "indexer.k_norm.weight": "F32", "ple_norm_conv.weight": "F32", "ple_norm_key.weight": "F32",
+    "ple_norm_query.weight": "F32", "ssm_a": "F32", "ssm_conv1d.weight": "F32", "ssm_dt.bias": "F32",
+    "ssm_norm.weight": "F32",
+}
+# Served from the GGUF when quantized in a type the engine has kernels for: the dense projections and the PLE key
+# (NativeDense and NativeHead: native_mmvq_supported), the token embedding (NativeEmbed: iq_supported).
+MMVQ_TYPES = {"Q4_0", "Q5_0", "Q8_0", "Q3_K", "Q4_K", "Q5_K", "Q6_K", "IQ4_NL", "IQ4_XS", "Q2_0", "IQ2_XXS",
+              "IQ2_XS", "IQ3_XXS", "IQ3_S", "IQ2_S", "IQ1_M"}
+EMBED_TYPES = {"IQ2_XXS", "IQ2_XS", "IQ3_XXS", "IQ4_NL", "IQ3_S", "IQ2_S", "IQ4_XS", "IQ1_S", "IQ1_M", "Q2_0", "Q3_K",
+               "Q4_K", "Q5_K", "Q6_K", "Q5_1", "Q8_0"}
+NATIVE = {"attn_qkv.weight", "attn_gate.weight", "ssm_out.weight", "attn_q.weight", "attn_k.weight",
+          "attn_v.weight", "attn_output.weight", "ffn_gate_shexp.weight", "ffn_up_shexp.weight",
+          "ffn_down_shexp.weight", "ple_key.weight", "output.weight"}
+PLE_TABLE = "per_layer_token_embd.weight"   # read from the GGUF by the engine (--ple-gguf)
+
+
+class Model:
+    """A model's shards (a split model's are found by name from the first), tensors looked up across them."""
+
+    def __init__(self, first: pathlib.Path):
+        m = re.fullmatch(r"(.*-)00001(-of-(\d{5})\.gguf)", first.name)
+        names = [f"{m.group(1)}{i:05d}{m.group(2)}" for i in range(1, int(m.group(3)) + 1)] if m else [first.name]
+        self.shards = []
+        self.where = {}
+        for name in names:
+            path = first.parent / name
+            if not path.exists():
+                raise FileNotFoundError(f"missing model shard {path}")
+            g = G.GGUFFile(path)
+            mm = np.memmap(path, dtype=np.uint8, mode="r")
+            for t in g.tensors:
+                if t.name in self.where:
+                    raise ValueError(f"tensor {t.name} is in two shards")
+                self.where[t.name] = (g, mm, t)
+            self.shards.append(g)
+
+    def tensors(self):
+        return [t for g in self.shards for t in g.tensors]
+
+    def get(self, name):
+        return self.where[name][2] if name in self.where else None
+
+    def bytes(self, t) -> np.ndarray:
+        g, mm, _ = self.where[t.name]
+        return mm[g.data_start + t.offset: g.data_start + t.offset + t.expected_bytes()]
+
+
+def suffix(name: str) -> str:
+    return re.sub(r"^blk\.\d+\.", "", name)
+
+
+def to_f32(raw: np.ndarray, type_name: str) -> np.ndarray:
+    if type_name == "F32":
+        return raw.view("<f4")
+    if type_name == "F16":
+        return raw.view("<f2").astype(np.float32)
+    if type_name == "BF16":
+        return (raw.view("<u2").astype(np.uint32) << 16).view(np.float32)
+    if type_name == "Q8_0":
+        b = raw.reshape(-1, 34)
+        return (b[:, 2:].view(np.int8).astype(np.float32) * b[:, :2].copy().view("<f2").astype(np.float32)).reshape(-1)
+    raise ValueError(f"no conversion from {type_name}")
+
+
+def encode(x: np.ndarray, form: str) -> bytes:
+    if form == "F32":
+        return x.astype("<f4").tobytes()
+    if form == "F16":
+        return x.astype("<f2").tobytes()
+    u = x.astype(np.float32).view(np.uint32)                   # BF16, round to nearest even
+    return ((u + 0x7FFF + ((u >> 16) & 1)) >> 16).astype("<u2").tobytes()
+
+
+def is_expert(name: str) -> bool:
+    return name.startswith("blk.") and name.endswith(("_exps.weight",))


 def read_index(path: pathlib.Path):
@@ -50,40 +145,52 @@ def read_index(path: pathlib.Path):
    return header, rows


-def tensor_bytes(mm, g, t) -> np.ndarray:
-    n = t.expected_bytes()
-    return mm[g.data_start + t.offset: g.data_start + t.offset + n]
+def native_row(t) -> list[str]:
+    ne0 = int(t.shape[0])
+    ne1 = int(t.shape[1]) if len(t.shape) > 1 else 0
+    return [t.name, "0", "0", "0", "0", "0", "0", str(ne0), str(ne1), "8", "0", "32"] + ["0"] * 7


-def is_expert(name: str) -> bool:
-    return name.startswith("blk.") and name.endswith(("_exps.weight",))
-
-
-def index_standalone(src, out, g, T, mm) -> int:
-    """Every non-expert tensor of shard 1: floats into dense.bin as stored, quantized ones native-only."""
-    rows, at = [], 0
-    served = 0
+def index_standalone(src, out, model: Model) -> int:
+    """Every non-expert tensor: quantized ones the engine serves natively as shape-only rows, the others into
+    dense.bin in their engine form."""
+    rows, at, served, converted = [], 0, 0, 0
    with open(out / "dense.bin", "wb") as fo:
-        for t in g.tensors:
-            if is_expert(t.name):
+        for t in model.tensors():
+            if is_expert(t.name) or t.name == PLE_TABLE:
                continue
            if len(t.shape) > 2:
                print("tensor %s has %d dimensions; the index holds two" % (t.name, len(t.shape)))
                return 1
+            sfx = suffix(t.name)
+            if t.name == "token_embd.weight" or (sfx in NATIVE and t.type_name not in FLOAT):
+                ok = EMBED_TYPES if t.name == "token_embd.weight" else MMVQ_TYPES
+                if t.type_name not in ok:
+                    print("%s is %s, a type the engine has no native kernels for" % (t.name, t.type_name))
+                    return 1
+                served += 1
+                rows.append(native_row(t))
+                continue
+            form = FORM.get(sfx)
+            if form is None:
+                print("tensor %s: no engine form known for it" % t.name)
+                return 1
+            raw = model.bytes(t)
+            if t.type_name == form:
+                data = raw.tobytes()
+            else:
+                data = encode(to_f32(raw, t.type_name), form)
+                converted += 1
            ne0 = int(t.shape[0])
            ne1 = int(t.shape[1]) if len(t.shape) > 1 else 0
-            if t.type_name in FLOAT:
-                raw = tensor_bytes(mm, g, t).tobytes()
-                kind = {"BF16": "4", "F16": "5", "F32": "2"}[t.type_name]
-                rows.append([t.name, "0", kind, str(at), str(len(raw)), "0", str(len(raw)), str(ne0), str(ne1),
-                             "0", "0", "1"] + ["0"] * 7)
-                fo.write(raw)
-                pad = (-len(raw)) % ALIGN
-                fo.write(b"\0" * pad)
-                at += len(raw) + pad
-            else:
-                served += 1
-                rows.append([t.name, "0", "0", "0", "0", "0", "0", str(ne0), str(ne1), "8", "0", "32"] + ["0"] * 7)
+            rows.append([t.name, "0", KIND[form], str(at), str(len(data)), "0", str(len(data)), str(ne0), str(ne1),
+                         "0", "0", "1"] + ["0"] * 7)
+            fo.write(data)
+            pad = (-len(data)) % ALIGN
+            fo.write(b"\0" * pad)
+            at += len(data) + pad
+    if converted:
+        print("dense.bin: %d tensors converted to their engine form" % converted)
    write_index(out, rows, src, served, 0)
    return 0

@@ -102,25 +209,23 @@ def write_index(out, rows, src, served, n_extra):
          % (len(rows), served, n_extra, at / 2**30))


-def index_from_base(a, src, base, out, g, T, mm) -> int:
+def index_from_base(src, base, out, model: Model) -> int:
    base_src = pathlib.Path(json.loads((base / "manifest.json").read_text(encoding="utf-8"))["source"]["shard1"])
    if not base_src.exists():
        print("cannot find the base pack's shard 1 from its manifest.json")
        return 1
-    bg = G.GGUFFile(base_src)
-    BT = {t.name: t for t in bg.tensors}
-    bmm = np.memmap(base_src, dtype=np.uint8, mode="r")
+    bm = Model(base_src)
    header, rows = read_index(base / "index.txt")
    new_rows, extra = [], []
    served = 0
    for name, f in rows.items():
-        t, bt = T.get(name), BT.get(name)
+        t, bt = model.get(name), bm.get(name)
        if t is None or bt is None:
            print("tensor %s missing from one of the models" % name)
            return 1
        if t.type_name in FLOAT and bt.type_name in FLOAT:
            if t.type_name != bt.type_name or t.shape != bt.shape or \
-                    not np.array_equal(tensor_bytes(mm, g, t), tensor_bytes(bmm, bg, bt)):
+                    not np.array_equal(model.bytes(t), bm.bytes(bt)):
                print("float tensor %s differs from the base model; this pack cannot reuse its dense.bin" % name)
                return 1
            new_rows.append(list(f))
@@ -130,7 +235,7 @@ def index_from_base(a, src, base, out, g, T, mm) -> int:
                return 1
            nbytes = t.expected_bytes()
            off = sum(len(b) + (-len(b)) % ALIGN for b in extra)
-            extra.append(tensor_bytes(mm, g, t).tobytes())
+            extra.append(model.bytes(t).tobytes())
            # file 3 = extra.bin, raw BF16 (index kind 4)
            new_rows.append([name, "3", "4", str(off), str(nbytes), "0", str(nbytes), f[7], f[8]] + ["0"] * 10)
        else:
@@ -154,7 +259,7 @@ def index_from_base(a, src, base, out, g, T, mm) -> int:

 def main() -> int:
    ap = argparse.ArgumentParser()
-    ap.add_argument("--gguf", required=True, help="the model's shard 1")
+    ap.add_argument("--gguf", required=True, help="the model's first shard (the others are found by name)")
    ap.add_argument("--base", help="optional: a Q2_0 canonical pack whose dense.bin holds the shared float tensors")
    ap.add_argument("--out", required=True)
    ap.add_argument("--skip-experts", action="store_true", help="rewrite the index only")
@@ -166,13 +271,8 @@ def main() -> int:
    out = pathlib.Path(a.out)
    out.mkdir(parents=True, exist_ok=True)

-    g = G.GGUFFile(src)
-    T = {t.name: t for t in g.tensors}
-    mm = np.memmap(src, dtype=np.uint8, mode="r")
-    if a.base:
-        rc = index_from_base(a, src, base, out, g, T, mm)
-    else:
-        rc = index_standalone(src, out, g, T, mm)
+    model = Model(src)
+    rc = index_from_base(src, base, out, model) if a.base else index_standalone(src, out, model)
    if rc:
        return rc
    if not (out / "tokenizer" / "vocab.json").exists():
@@ -180,10 +280,11 @@ def main() -> int:
                       check=True)   # writes <out>/tokenizer/

    # ---- the experts
-    n_layers = 1 + max(int(n.split(".")[1]) for n in T if n.startswith("blk.") and n.endswith("_exps.weight"))
+    names = [t.name for t in model.tensors()]
+    n_layers = 1 + max(int(n.split(".")[1]) for n in names if n.startswith("blk.") and n.endswith("_exps.weight"))
    layout, offset = [], 0
    for l in range(n_layers):
-        ts = [T["blk.%d.ffn_%s_exps.weight" % (l, r)] for r in ROLES]
+        ts = [model.get("blk.%d.ffn_%s_exps.weight" % (l, r)) for r in ROLES]
        per = [t.expected_bytes() // N_EXPERT for t in ts]
        if per[0] != per[1] or ts[0].type_name != ts[1].type_name:
            print("layer %d: gate and up differ in type" % l)
@@ -192,10 +293,11 @@ def main() -> int:
        layout.append((l, ts[0].type_id, ts[2].type_id, offset, blob, ts))
        offset += blob * N_EXPERT
    with open(out / "native_experts.txt", "w", encoding="utf-8", newline="\n") as fo:
-        fo.write("# strata native experts v2: layer gu_type d_type offset blob_bytes gate_off up_off down_off "
-                 "(n_expert %d, total %d; the last three are absolute offsets in %s)\n" % (N_EXPERT, offset, src.name))
+        fo.write("# strata native experts v3: layer gu_type d_type offset blob_bytes (n_expert %d, total %d; the "
+                 "engine takes the experts from %s by tensor name, or from experts.bin)\n"
+                 % (N_EXPERT, offset, src.name))
        for l, gt, dt, off, blob, ts in layout:
-            fo.write("%d %d %d %d %d %d %d %d\n" % (l, gt, dt, off, blob, *[g.data_start + t.offset for t in ts]))
+            fo.write("%d %d %d %d %d\n" % (l, gt, dt, off, blob))
    if a.skip_experts or not a.experts_bin:
        if (out / "experts.bin").exists() and not a.experts_bin:
            print("note: %s/experts.bin exists; the engine reads it instead of the GGUF" % out)
@@ -206,7 +308,7 @@ def main() -> int:
        return 0
    with open(path, "wb") as fo:
        for l, gt, dt, off, blob, ts in layout:
-            parts = [tensor_bytes(mm, g, t).reshape(N_EXPERT, -1) for t in ts]
+            parts = [model.bytes(t).reshape(N_EXPERT, -1) for t in ts]
            chunk = np.concatenate(parts, axis=1)          # (512, blob): gate | up | down per expert
            assert chunk.shape == (N_EXPERT, blob)
            fo.write(chunk.tobytes())