Add files using upload-large-folder tool

MTP-VERIFY-REPORT.json

@@ -0,0 +1,17 @@
+{
+  "status": "PASS",
+  "layer": 45,
+  "experts": 288,
+  "sidecars": 4,
+  "expert_stream_mismatches": 0,
+  "codec_closure": "PASS: all 864 projection encodes",
+  "coupled_component_validation": "PASS",
+  "routed_sidecar_bytes": 3853990880,
+  "mean_tw_nmse": {
+    "gate_tw_nmse": 0.005752041559620415,
+    "up_tw_nmse": 0.005751395422463146,
+    "down_tw_nmse": 0.0061286417262067724
+  },
+  "scale_policy": "signed-unit coupled vectors; CPU torch seed 530045; no prior r27 MTP scales",
+  "native_execution_tested": false
+}

VERIFY-REPORT.json

@@ -1,56 +1,74 @@
 {
- "checkpoint": "/root/cv2data/out/ckpt-B",
- "runtime_overlay_and_kernel_header_checks": "PASS (168 files)",
- "expert_stream_mismatches": {
-  "3": 0,
-  "4": 0,
-  "5": 0,
-  "6": 0,
-  "7": 0,
-  "8": 0,
-  "9": 0,
-  "10": 0,
-  "11": 0,
-  "12": 0,
-  "13": 0,
-  "14": 0,
-  "15": 0,
-  "16": 0,
-  "17": 0,
-  "18": 0,
-  "19": 0,
-  "20": 0,
-  "21": 0,
-  "22": 0,
-  "23": 0,
-  "24": 0,
-  "25": 0,
-  "26": 0,
-  "27": 0,
-  "28": 0,
-  "29": 0,
-  "30": 0,
-  "31": 0,
-  "32": 0,
-  "33": 0,
-  "34": 0,
-  "35": 0,
-  "36": 0,
-  "37": 0,
-  "38": 0,
-  "39": 0,
-  "40": 0,
-  "41": 0,
-  "42": 0,
-  "43": 0,
-  "44": 0
- },
- "routed_sidecar_bytes": 161867586048,
- "tp2_routed_bytes_per_gpu": 80933793024.0,
- "allocation_counts": {
-  "3": 0,
-  "4": 42,
-  "5": 0
- },
- "status": "PASS"
+  "checkpoint": "/root/cv2data/out/ckpt-B",
+  "runtime_overlay_and_kernel_header_checks": "PASS: 168 main files plus 4 MTP files",
+  "expert_stream_mismatches": {
+    "3": 0,
+    "4": 0,
+    "5": 0,
+    "6": 0,
+    "7": 0,
+    "8": 0,
+    "9": 0,
+    "10": 0,
+    "11": 0,
+    "12": 0,
+    "13": 0,
+    "14": 0,
+    "15": 0,
+    "16": 0,
+    "17": 0,
+    "18": 0,
+    "19": 0,
+    "20": 0,
+    "21": 0,
+    "22": 0,
+    "23": 0,
+    "24": 0,
+    "25": 0,
+    "26": 0,
+    "27": 0,
+    "28": 0,
+    "29": 0,
+    "30": 0,
+    "31": 0,
+    "32": 0,
+    "33": 0,
+    "34": 0,
+    "35": 0,
+    "36": 0,
+    "37": 0,
+    "38": 0,
+    "39": 0,
+    "40": 0,
+    "41": 0,
+    "42": 0,
+    "43": 0,
+    "44": 0,
+    "45": 0
+  },
+  "routed_sidecar_bytes": 165721576928,
+  "tp2_routed_bytes_per_gpu": 82860788464.0,
+  "allocation_counts": {
+    "3": 0,
+    "4": 43,
+    "5": 0
+  },
+  "status": "PASS",
+  "mtp_verification": {
+    "status": "PASS",
+    "layer": 45,
+    "experts": 288,
+    "sidecars": 4,
+    "expert_stream_mismatches": 0,
+    "codec_closure": "PASS: all 864 projection encodes",
+    "coupled_component_validation": "PASS",
+    "routed_sidecar_bytes": 3853990880,
+    "mean_tw_nmse": {
+      "gate_tw_nmse": 0.005752041559620415,
+      "up_tw_nmse": 0.005751395422463146,
+      "down_tw_nmse": 0.0061286417262067724
+    },
+    "scale_policy": "signed-unit coupled vectors; CPU torch seed 530045; no prior r27 MTP scales",
+    "native_execution_tested": false
+  }
 }

carrier/model-mtp-nonexpert.safetensors

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:31fc8804fce4cb1e0207d9fb13c19a892d85b50a7f769e763cb8d9e5b6a2e6ba
+size 369674104

evidence/build-mtp-checkpoint.py

@@ -0,0 +1,144 @@
+"""Assemble and verify MTP45 sidecars and retain only its unchanged carrier tensors."""
+import hashlib
+import json
+import math
+import os
+import shutil
+import time
+from pathlib import Path
+
+import torch
+from safetensors import safe_open
+from safetensors.torch import load_file, save_file
+
+from build_checkpoint import build_rank, sha256_file, tensor_sha256, CODEC
+from p8_coupled_scales import SCALE_NAMES, validate_coupled_component
+
+out = Path(os.environ["CV2_OUT"])
+root = Path(os.environ["CV2_ROOT"])
+while not (out / "status/mtp-encoded.ready").exists():
+    time.sleep(10)
+inputs = out / "mtp-inputs"
+input_info = json.loads((inputs / "inputs.json").read_text())
+core_sha = sha256_file(root / "code/codecv2/encode_layer_v2.py")
+wrapper_sha = sha256_file(root / "ops/encode-mtp-v2.py")
+chunks, metrics = {}, []
+for path in sorted((out / "mtp-chunks").glob("*.safetensors")):
+    with safe_open(str(path), framework="pt", device="cpu") as src:
+        meta = src.metadata()
+        assert meta["schema"] == "trellismx-codec-v2-mtp-chunk.v1"
+        assert meta["layer"] == "45" and meta["bits"] == "4"
+        assert meta["encoder_sha256"] == core_sha and meta["mtp_wrapper_sha256"] == wrapper_sha
+        assert meta["mtp_inputs_sha256"] == sha256_file(inputs / "inputs.json")
+        assert float(meta["beta"]) == 0.5 and float(meta["percdamp"]) == 0.3
+        start, end = map(int, meta["expert_range"].split(":"))
+        for expert in range(start, end):
+            assert expert not in chunks
+            chunks[expert] = {p: (src.get_tensor(f"{expert}.{p}.trellis"), src.get_tensor(f"{expert}.{p}.scale_ue8m0"))
+                              for p in ("gate", "up", "down")}
+    metrics.extend(json.loads(path.with_suffix(".metrics.json").read_text())["rows"])
+assert set(chunks) == set(range(288))
+assert {r["expert"] for r in metrics} == set(range(288)) and len(metrics) == 288
+assert all(math.isfinite(v) for row in metrics for v in row.values())
+dest = out / "mtp-sidecars"
+(dest / "sidecars").mkdir(parents=True, exist_ok=True)
+(dest / "design").mkdir(exist_ok=True)
+design = {"schema": "trellismx-codecv2-mtp-design.v1", "layer": 45, "bits": 4,
+          "encoder_core_sha256": core_sha, "mtp_wrapper_sha256": wrapper_sha,
+          "beta": 0.5, "percdamp": 0.3, "coupled_sign_draw": 0,
+          "scale_policy": input_info["scale_policy"], "scale_sha256": input_info["scale_sha256"],
+          "capture_sha256": input_info["source_capture_sha256"],
+          "calibration": "64 fit windows; 2047 valid MTP rows each; all routed plus every fourth all-token row",
+          "main_model_scales_unchanged": True}
+design_path = dest / "design/design-2.json"
+design_path.write_text(json.dumps(design, indent=2) + "\n")
+design_sha = sha256_file(design_path)
+scale = load_file(str(inputs / "coupled-scales.safetensors"))
+transform_sha = sha256_file(out / "ckpt-B/design/transform.json")
+records, ranks = [], []
+for rank in range(4):
+    template = out / "ckpt-B/sidecars" / f"p8-layer-044-tp4-rank-{rank}.safetensors"
+    tensors, meta = build_rank(template, chunks, rank, 4, design_sha)
+    sl = slice(rank * 512, (rank + 1) * 512)
+    tensors["gate_up_suh_fp16"] = scale["gate_up_suh"]
+    tensors["down_svh_fp16"] = scale["down_svh"]
+    tensors["intermediate_scales_fp16"] = torch.cat([scale[key][:, sl] for key in ("gate_svh", "up_svh", "down_suh")], dim=1).contiguous()
+    meta.pop("exl3_scale_source_sha256", None)
+    meta.update(layer="45", mtp_scale_source_sha256=input_info["scale_sha256"],
+                mtp_scale_policy=input_info["scale_policy"], mtp_capture_sha256=input_info["source_capture_sha256"])
+    for name in SCALE_NAMES:
+        meta["sha256_" + name] = tensor_sha256(tensors[name])
+    validate_coupled_component(meta, {name: tensors[name] for name in SCALE_NAMES},
+        layer=45, rank=rank, experts=288, hidden=4096, intermediate=512,
+        expected_transform_sha256=transform_sha)
+    name = f"sidecars/p8-layer-045-tp4-rank-{rank}.safetensors"
+    path = dest / name
+    save_file(tensors, str(path), metadata=meta)
+    records.append({"layer": 45, "rank": rank, "bits": 4, "path": name,
+                    "bytes": path.stat().st_size, "sha256": sha256_file(path), "source_design_sha256": design_sha})
+    # Read the actual saved artifact for the full stream reassembly check.
+    with safe_open(str(path), framework="pt", device="cpu") as src:
+        ranks.append({key: src.get_tensor(key) for key in CODEC})
+    print(json.dumps({"saved": name, "bytes": path.stat().st_size}), flush=True)
+mismatches = 0
+for expert in range(288):
+    actual = {
+        "gate": (torch.cat([t["w13_trellis"][0, expert] for t in ranks], 1),
+                 torch.cat([t["w13_scale_ue8m0"][expert, 512:] for t in ranks], 0)),
+        "up": (torch.cat([t["w13_trellis"][1, expert] for t in ranks], 1),
+               torch.cat([t["w13_scale_ue8m0"][expert, :512] for t in ranks], 0)),
+        "down": (torch.cat([t["w2_trellis"][expert] for t in ranks], 0),
+                 torch.cat([t["w2_scale_ue8m0"][expert] for t in ranks], 1))}
+    mismatches += not all(torch.equal(a, b) for key in actual for a, b in zip(actual[key], chunks[expert][key]))
+assert mismatches == 0
+report = {"status": "PASS", "layer": 45, "experts": 288, "sidecars": 4,
+          "expert_stream_mismatches": mismatches, "codec_closure": "PASS: all 864 projection encodes",
+          "coupled_component_validation": "PASS", "routed_sidecar_bytes": sum(r["bytes"] for r in records),
+          "mean_tw_nmse": {key: sum(r[key] for r in metrics) / 288 for key in ("gate_tw_nmse", "up_tw_nmse", "down_tw_nmse")},
+          "scale_policy": input_info["scale_policy"], "native_execution_tested": False}
+(dest / "MTP-VERIFY-REPORT.json").write_text(json.dumps(report, indent=2) + "\n")
+(dest / "mtp-manifest.json").write_text(json.dumps({"files": records, "design_sha256": design_sha}, indent=2) + "\n")
+
+# Remove only replaced MTP expert tensors; retain all 25 other MTP tensors bit for bit.
+old_carrier = out / "native-carrier"
+carrier = out / "native-carrier-mtp"
+carrier.mkdir(exist_ok=True)
+base_receipt = json.loads((out / "native-carrier-receipt.json").read_text())
+kept_tensors = {}
+for record in base_receipt["files"]:
+    name = record["path"]
+    if name.startswith("model-mtp-"):
+        with safe_open(str(old_carrier / name), framework="pt", device="cpu") as src:
+            for key in src.keys():
+                if ".layers.45.mlp.experts." not in key:
+                    assert key not in kept_tensors
+                    kept_tensors[key] = src.get_tensor(key)
+    elif name != "model.safetensors.index.json":
+        if not (carrier / name).exists():
+            os.link(old_carrier / name, carrier / name)
+assert len(kept_tensors) == 25
+mtp_file = "model-mtp-nonexpert.safetensors"
+save_file(kept_tensors, str(carrier / mtp_file))
+with safe_open(str(carrier / mtp_file), framework="pt", device="cpu") as check:
+    assert set(check.keys()) == set(kept_tensors)
+    assert all(torch.equal(check.get_tensor(key), value) for key, value in kept_tensors.items())
+index = json.loads((old_carrier / "model.safetensors.index.json").read_text())
+weight_map = {k: v for k, v in index["weight_map"].items() if ".layers.45.mlp.experts." not in k}
+for key in kept_tensors:
+    weight_map[key] = mtp_file
+index["weight_map"] = weight_map
+index["metadata"]["total_size"] -= 7247757312 + 226492416
+(carrier / "model.safetensors.index.json").write_text(json.dumps(index, indent=2) + "\n")
+known = {r["path"]: r for r in base_receipt["files"]}
+new_records = []
+for path in sorted(carrier.iterdir()):
+    assert path.is_file()
+    digest = known[path.name]["sha256"] if path.name in known and path.name != "model.safetensors.index.json" else sha256_file(path)
+    new_records.append({"path": path.name, "bytes": path.stat().st_size, "sha256": digest})
+receipt = {**base_receipt, "files": new_records, "bytes": sum(r["bytes"] for r in new_records),
+           "indexed_tensors": len(weight_map), "mtp_expert_tensors_replaced": 1728,
+           "mtp_nonexpert_tensors_preserved": 25, "mtp_nonexpert_tensor_equality": "PASS",
+           "derivation": "MTP45 FP8 experts and scales replaced by Trellis sidecars; other tensors unchanged"}
+(out / "native-carrier-mtp-receipt.json").write_text(json.dumps(receipt, indent=2) + "\n")
+(out / "status/mtp.ready").write_text(design_sha + "\n")
+print(json.dumps(report), flush=True)

evidence/encode-mtp-v2.py

@@ -0,0 +1,91 @@
+"""Encode MTP45 with the supplied codec-v2 math and its 2047-row fit captures."""
+import argparse
+import hashlib
+import json
+import os
+import time
+from pathlib import Path
+
+import numpy as np
+import torch
+from safetensors.torch import load_file, save_file
+
+import encode_layer_v2 as cv2
+from glm53_nvfp4.p8_coupled_scale import CoupledScaleSet, COUPLED_SIGN_DRAW, encode_coupled_scale_weights, coupled_input_carrier, coupled_middle_carrier
+from glm53_nvfp4.shard_index import IndexedCheckpoint
+
+ap = argparse.ArgumentParser()
+ap.add_argument("--expert-start", type=int, required=True)
+ap.add_argument("--expert-end", type=int, required=True)
+ap.add_argument("--out", type=Path, required=True)
+args = ap.parse_args()
+assert 0 <= args.expert_start < args.expert_end <= 288
+assert not args.out.exists()
+inputs = Path(os.environ["CV2_OUT"]) / "mtp-inputs"
+info = json.loads((inputs / "inputs.json").read_text())
+assert info["rows"] == 64 * 2047
+hidden = np.memmap(inputs / "hidden.bf16.bin", mode="r", dtype="<u2", shape=(info["rows"], 4096))
+ids = np.memmap(inputs / "topk_ids.u16le.bin", mode="r", dtype="<u2", shape=(info["rows"], 8))
+routes = np.memmap(inputs / "topk_weights.f32le.bin", mode="r", dtype="<f4", shape=(info["rows"], 8))
+scales = load_file(str(inputs / "coupled-scales.safetensors"))
+scales_all = CoupledScaleSet(**scales, source_path=inputs, source_sha256=info["scale_sha256"], source_metadata={}, tensor_hashes={})
+scales_all.validate()
+device = torch.device("cuda:0")
+torch.cuda.set_device(device)
+every = np.asarray([w * 2047 + t for w in range(64) for t in range(0, 2047, 4)])
+hid_every = torch.from_numpy(np.array(hidden[every], copy=True)).view(torch.bfloat16)
+ones_every = torch.ones(len(every))
+source_path = Path(os.environ["CV2_BF16"])
+source = IndexedCheckpoint(source_path, source_path / "model.safetensors.index.json")
+payload, metrics, h_all_in = {}, [], None
+started = time.time()
+for expert in range(args.expert_start, args.expert_end):
+    t0 = time.time()
+    base = source.expert_prefix(45, expert)
+    orig = {p: source.get(f"{base}{p}.weight").to(device).float() for p in cv2.PROJ}
+    sc = CoupledScaleSet(gate_up_suh=scales_all.gate_up_suh.to(device),
+        down_svh=scales_all.down_svh.to(device),
+        gate_svh=scales_all.gate_svh[expert:expert + 1].to(device),
+        up_svh=scales_all.up_svh[expert:expert + 1].to(device),
+        down_suh=scales_all.down_suh[expert:expert + 1].to(device),
+        source_path=inputs, source_sha256=info["scale_sha256"], source_metadata={}, tensor_hashes={})
+    tw = encode_coupled_scale_weights(orig["gate_proj"], orig["up_proj"], orig["down_proj"], sc,
+                                       expert=0, intermediate_draw=COUPLED_SIGN_DRAW)
+    carrier = lambda h: coupled_input_carrier(h, sc.gate_up_suh, quantize=True)
+    row_ids, slots = np.nonzero(ids == expert)
+    assert row_ids.size, f"No MTP fit routes for expert {expert}"
+    hid_r = torch.from_numpy(np.array(hidden[row_ids], copy=True)).view(torch.bfloat16)
+    rt_r = torch.from_numpy(np.array(routes[row_ids, slots], copy=True))
+    if h_all_in is None:
+        h_all_in = cv2.hessian_chunks(hid_every, ones_every, carrier, device)
+    h_in = cv2.mix(cv2.hessian_chunks(hid_r, rt_r, carrier, device), h_all_in, 0.5)
+    g_rec, g_tr, g_sc = cv2.encode_v2(tw[0], h_in, 4, 0.3)
+    u_rec, u_tr, u_sc = cv2.encode_v2(tw[1], h_in, 4, 0.3)
+    mid = lambda h: coupled_middle_carrier(carrier(h), g_rec, u_rec, sc, expert=0,
+                                           intermediate_draw=COUPLED_SIGN_DRAW, quantize=True)
+    h_dn = cv2.mix(cv2.hessian_chunks(hid_r, rt_r, mid, device),
+                   cv2.hessian_chunks(hid_every, ones_every, mid, device), 0.5)
+    d_rec, d_tr, d_sc = cv2.encode_v2(tw[2], h_dn, 4, 0.3)
+    row = {"expert": expert, "routed_tokens": len(row_ids), "seconds": round(time.time() - t0, 2)}
+    for name, rec, tr, codes, target in (("gate", g_rec, g_tr, g_sc, tw[0]),
+            ("up", u_rec, u_tr, u_sc, tw[1]), ("down", d_rec, d_tr, d_sc, tw[2])):
+        payload[f"{expert}.{name}.trellis"] = tr
+        payload[f"{expert}.{name}.scale_ue8m0"] = codes
+        row[f"{name}_tw_nmse"] = float((rec - target).double().square().sum() / target.double().square().sum())
+    assert all(np.isfinite(v) for v in row.values())
+    metrics.append(row)
+    print(json.dumps(row), flush=True)
+    del orig, tw, h_in, h_dn, g_rec, u_rec, d_rec, hid_r, rt_r
+    torch.cuda.empty_cache()
+meta = {"schema": "trellismx-codec-v2-mtp-chunk.v1", "layer": "45", "bits": "4",
+        "expert_range": f"{args.expert_start}:{args.expert_end}", "beta": "0.5", "percdamp": "0.3",
+        "encoder_sha256": hashlib.sha256(Path(cv2.__file__).read_bytes()).hexdigest(),
+        "mtp_wrapper_sha256": hashlib.sha256(Path(__file__).read_bytes()).hexdigest(),
+        "mtp_inputs_sha256": hashlib.sha256((inputs / "inputs.json").read_bytes()).hexdigest(),
+        "scale_policy": info["scale_policy"]}
+args.out.parent.mkdir(parents=True, exist_ok=True)
+tmp = args.out.with_suffix(".partial")
+save_file(payload, str(tmp), metadata=meta)
+tmp.rename(args.out)
+args.out.with_suffix(".metrics.json").write_text(json.dumps({"meta": meta, "rows": metrics, "seconds": time.time() - started}, indent=2) + "\n")
+print(json.dumps({"done": str(args.out), "seconds": time.time() - started}), flush=True)

evidence/mtp-inputs.json

@@ -0,0 +1,717 @@
+{
+  "schema": "trellismx-codecv2-mtp-inputs.v1",
+  "layer": 45,
+  "capture_repo": "brandonmusic/GLM-5.3-Flash-BF16-Teacher-Logits",
+  "capture_revision": "95f4fdd94bf29989db2e0d1054e4931f55edb6aa",
+  "source_capture_sha256": "50afb51ecb3519599bf8722e7f2fe92e3e0edc96d3e3ba605ed59246055f8ee8",
+  "windows": [
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 6144,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "208eadbaf0f326e410557519bd28398a6cb44027a143fa7d5eb4e104e492c59f",
+      "window_id": "fit-0003",
+      "window_index": 3
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 14336,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "b4b12b709b0643763ea88f28a7159fb0d6f67ec403a769d82888097330c588fe",
+      "window_id": "fit-0007",
+      "window_index": 7
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 22528,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "03e62368844238497ffcf7a12bd9baa5bd36480ec68d2f8563c9180fd275560f",
+      "window_id": "fit-0011",
+      "window_index": 11
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 30720,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "d35fae888ad057534d85e66c7eec53d189eb31f4312ed9f84009501f20584619",
+      "window_id": "fit-0015",
+      "window_index": 15
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 38912,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "ff92c88a67993287542987983aa46d9e9a0ab636ff3ff453f0db487cd3654181",
+      "window_id": "fit-0019",
+      "window_index": 19
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 40960,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "c0f29007388fc4a03964260f8d9d5bc5e24239db854979e215866eeb4c8db3e9",
+      "window_id": "fit-0020",
+      "window_index": 20
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 45056,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "0ed0081f3ee1f0117a8b716fb530cbfa6b014d2a550e6a4a32f0e4c403559edd",
+      "window_id": "fit-0022",
+      "window_index": 22
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 47104,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "cb81685f194c02bd698d50ed0f5ec197a28dcfdca6d9a37c1af44111d9b5e438",
+      "window_id": "fit-0023",
+      "window_index": 23
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 55296,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "2a5e68ea816bcd1ae71eb65a2c2963ad693139fb9b18853cc395772310584346",
+      "window_id": "fit-0027",
+      "window_index": 27
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 63488,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "6bcd57c39f9f206255f13684ce2c00367879867f68b405589edd41b5678ac249",
+      "window_id": "fit-0031",
+      "window_index": 31
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 71680,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "e8866385831f1d5bfccba77298e19ac414b8b90dd468ef138d979f704675ce5f",
+      "window_id": "fit-0035",
+      "window_index": 35
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 79872,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "d2a9535df825df10239ecd7212ec5e0cca0e216e5303c6d7ec77cef63680cc06",
+      "window_id": "fit-0039",
+      "window_index": 39
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 88064,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "4388a3c18f9cdc932620967f6fda46146d90f722b8263482354eccdd16177c85",
+      "window_id": "fit-0043",
+      "window_index": 43
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 96256,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "0432ef02690e147e5b6d03c38f0207010e830c1344823223141b5ba49e771826",
+      "window_id": "fit-0047",
+      "window_index": 47
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 104448,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "5dd6cf9058c192340985bdc02d379b102c194c7e4da43ff209cf5de019b25b8c",
+      "window_id": "fit-0051",
+      "window_index": 51
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 106496,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "83b96172a361a5f1a8eac603c4ef5b1e96023cdd07faf388479d0147a8c289cb",
+      "window_id": "fit-0052",
+      "window_index": 52
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 112640,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "7472e1d9f1fdeb708fb933d8bf51682e4d498f4cd96b001c86e166936da4c122",
+      "window_id": "fit-0055",
+      "window_index": 55
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 120832,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "6455be9da5b596f57b219d3a01f25cc7e38d83f09738b4a041f6ddb72d5a909c",
+      "window_id": "fit-0059",
+      "window_index": 59
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis4_reasoning_termination-1",
+      "domain": "axis4_reasoning_termination",
+      "main_terminal_offset": 129024,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "f838036e2eb0726a528ea1e5252e58d3340fd57f0a8f655bfb16eba68b74dc3b",
+      "window_id": "fit-0063",
+      "window_index": 63
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 145408,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "2b9498a4c264be4b3336e94d572e22fedb24fd745537ab519fe52616c7fa8369",
+      "window_id": "fit-0071",
+      "window_index": 71
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 161792,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "51ba8b2845d334157b51261bd6f7f8144bda32f0b53ec082356f9f3bffbca0f9",
+      "window_id": "fit-0079",
+      "window_index": 79
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 190464,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "7675e123eef3e202d31be6af29c815254ae601a85119eab63f124b476af9cba2",
+      "window_id": "fit-0093",
+      "window_index": 93
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 235520,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "52c57e4c50a2f724e6b0bb02ebeff3c850a27f0b3eb95db9b33dfd2ecc36fdc6",
+      "window_id": "fit-0115",
+      "window_index": 115
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 237568,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "fa0c3db2b83b1236825127432a325fd888cf8f5233f1b0f9aa88e3d6c46dd816",
+      "window_id": "fit-0116",
+      "window_index": 116
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 249856,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "1767d9eceb5d6d223d1a6222dd7e71fec1dd775d8eab9b8b56b39064e7e41957",
+      "window_id": "fit-0122",
+      "window_index": 122
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 262144,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "5af6ec5c7649361f07f5217411a5ac85b9c5b4b9740967c0b8d99aba570866fc",
+      "window_id": "fit-0128",
+      "window_index": 128
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 268288,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "5fb1a5379d4fe994a1e86a2e1cf44aa56c52b7e02ad43a05df7c40bf43418e4f",
+      "window_id": "fit-0131",
+      "window_index": 131
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 274432,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "284ff8b2bdae3627c57a9da5b3e97bf54068681d41c65b9e27659aef73763581",
+      "window_id": "fit-0134",
+      "window_index": 134
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 282624,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "f8c769d67636da9dfca5d73e57292a34830cfeb60da4b67cba5f2df4b083cbe9",
+      "window_id": "fit-0138",
+      "window_index": 138
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 284672,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "639984e8d3351390d01d2f934cbc7d4635a2b281692326080ca0bee7fa82e365",
+      "window_id": "fit-0139",
+      "window_index": 139
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 292864,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "d98dc54bc9af44ae5f06bffa644c90d68ded7f231ee86daef18228b9104fd018",
+      "window_id": "fit-0143",
+      "window_index": 143
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 301056,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "a14a85dbd0f53c19dba5c9f952fcecc0bb9ab2ee419b3138278878fbbc46bede",
+      "window_id": "fit-0147",
+      "window_index": 147
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 317440,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "f18685915b8f8b2c8dbc7c411471d5c6e56bfdd9c3158241ad9cc7adf9bf1a85",
+      "window_id": "fit-0155",
+      "window_index": 155
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 319488,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "5e9bf81a13d603aeee09856b17cd8085475c8252af6f4376565886cf7c8c824c",
+      "window_id": "fit-0156",
+      "window_index": 156
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 333824,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "d3274fddf9feaff8a8460e7cb8dcec6de21d7ee34d8df54c8a14784995bc93ff",
+      "window_id": "fit-0163",
+      "window_index": 163
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 352256,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "2f02a7a1854c920b37b555302659fb206af3227ad18c97f98f9d71487cf140f3",
+      "window_id": "fit-0172",
+      "window_index": 172
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 354304,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "b90385699a4e5ad8511be2f979d8b1eefb4f51f15009d0a9cfea504823f7e804",
+      "window_id": "fit-0173",
+      "window_index": 173
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 399360,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "e94b7adcc6c9b3ea370d72b3a93e0e587f3f15d5efa7a51dc5dbb851ef008d0b",
+      "window_id": "fit-0195",
+      "window_index": 195
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 403456,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "34364d64711a7325d1c600a90cb67e75cab6e57a5659cee1f186bdad4f92e9d3",
+      "window_id": "fit-0197",
+      "window_index": 197
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 413696,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "11bc2a5757f9ef787e1f21ef06fc6d13c6a174a12486dc8fbb43194775120123",
+      "window_id": "fit-0202",
+      "window_index": 202
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 421888,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "dd9e2596cf1f732177b22cc3abbc8d14df675f7478cb7273b8b02522516d8f59",
+      "window_id": "fit-0206",
+      "window_index": 206
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 430080,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "d7508373c2c195b1dff5dc0b35fb171e00348cd735296383336fa461cd678bdc",
+      "window_id": "fit-0210",
+      "window_index": 210
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 440320,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "81fc07989fe68186d98c90b7807d07520d83c8070821307ad2d69c9d1d90b120",
+      "window_id": "fit-0215",
+      "window_index": 215
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 458752,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "95c807d7a5aad5ea11c2e22591930669aaca061412d41754cdde15a6a4546900",
+      "window_id": "fit-0224",
+      "window_index": 224
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 464896,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "4a2158f357933bf460001d22ee4b6fcdd3609f03e0b417d8e04c2ce049dadf51",
+      "window_id": "fit-0227",
+      "window_index": 227
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 466944,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "1ef5b044392253389dd12a22c1978e706ab6584dbf301d6253d1099a742bbefe",
+      "window_id": "fit-0228",
+      "window_index": 228
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 485376,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "cb88cfc15311ce056f24e24a627de2fbc91f1eea2ced97785029d5ef29d1e02e",
+      "window_id": "fit-0237",
+      "window_index": 237
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 505856,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "a2a532c83873a63fbd1ecd3dc072059389886f2eb24fbe14ecf2f36f7e016f87",
+      "window_id": "fit-0247",
+      "window_index": 247
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 534528,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "17037b1512bacfd160cff6d58a099ae398cbeed5bc171ab1c8e0f0af8933b19e",
+      "window_id": "fit-0261",
+      "window_index": 261
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 542720,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "c64237273b209a945c7c818ebd1d422cb2ab50e4ebd7ef4b8401b242a9029e7b",
+      "window_id": "fit-0265",
+      "window_index": 265
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 563200,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "1e994c189fbd3d369abf496db5ff48750f71b252f645b09462230218985abdc2",
+      "window_id": "fit-0275",
+      "window_index": 275
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 577536,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "a3bcd38c2780d32246f3b8b630bad422b4dd444572d648ca21e720cbd4c28128",
+      "window_id": "fit-0282",
+      "window_index": 282
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 616448,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "a8fe54b5fc7f4fa5936b330d93d76da5dfdacd9712ef7a6f477600e910a2125b",
+      "window_id": "fit-0301",
+      "window_index": 301
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 647168,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "0060f819a176ec0abfa02ad323c9015afd7f5d6d6e06c99d819f16cb8f7a8a25",
+      "window_id": "fit-0316",
+      "window_index": 316
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis2_legal-1",
+      "domain": "axis2_legal",
+      "main_terminal_offset": 649216,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "f728ace6743429f911cdec7051a6dff7db0b0f5f8f915b42767fac4c95f8d849",
+      "window_id": "fit-0317",
+      "window_index": 317
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 677888,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "429eb9411d4ec92a26abf72cf302ef72d94f2df9c9011be2e9a91fcaaec7bece",
+      "window_id": "fit-0331",
+      "window_index": 331
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 708608,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "5655a824cde5dbc1600eee1f27b4b7cd0494dd1826139e93966984e62b46d789",
+      "window_id": "fit-0346",
+      "window_index": 346
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 712704,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "49baabdd793585f57ed18f17f025b09e16bd8eb26a1ab826e2072bb140b80026",
+      "window_id": "fit-0348",
+      "window_index": 348
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 737280,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "023afe71006b811629b43bb3d6d3044d120726ed525723961d5ff083ff4c32bc",
+      "window_id": "fit-0360",
+      "window_index": 360
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 739328,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "44e340f3fcb5a830cbccc7683a48b4691f3709bbaa9bff7fab2c34fdf03c3cf1",
+      "window_id": "fit-0361",
+      "window_index": 361
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 747520,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "eb638a132cf39bd4e4d14d6a976a1e0d47646ca1da63218c95cdcd66f24ddd9b",
+      "window_id": "fit-0365",
+      "window_index": 365
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 751616,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "59f18d04f36791b4dbb1a158692a0c6b2b53bf2c7b9ae63a51c3183270688225",
+      "window_id": "fit-0367",
+      "window_index": 367
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis1_general-1",
+      "domain": "axis1_general",
+      "main_terminal_offset": 765952,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "ea0b1e6680826d54e456803cae763837712748fc73f5477bb414ca708fea55aa",
+      "window_id": "fit-0374",
+      "window_index": 374
+    },
+    {
+      "attention_mask_sha256": "3f9144983ff37e8e104c28604ca9fc99957529a5c1e1a17904cc38f64b472a86",
+      "document_id": "reap-recall-packed-axis3_code_agentic-4",
+      "domain": "axis3_code_agentic",
+      "main_terminal_offset": 772096,
+      "role": "fit",
+      "rows": 2047,
+      "token_ids_sha256": "75ff1700624bc9c4e5251f3e8afe5b3fc6293199739cceaaca79de0c1ac6b175",
+      "window_id": "fit-0377",
+      "window_index": 377
+    }
+  ],
+  "rows_per_window": 2047,
+  "rows": 131008,
+  "scale_policy": "signed-unit coupled vectors; CPU torch seed 530045; no prior r27 MTP scales",
+  "scale_sha256": "ad170baa8259a6af3f234f154c0b9cca188b8d1f7cf6de0c89788bccbf66827b"
+}

evidence/native-carrier-mtp-receipt.json

@@ -0,0 +1,112 @@
+{
+  "repo": "local-inference-lab/GLM-5.3-Flash-NVFP4",
+  "revision": "520de24eabf507659eaef7c70f14fd584527facc",
+  "source_verification": "PASS",
+  "index_header_coverage": "PASS",
+  "indexed_tensors": 37906,
+  "omitted_replaced_routed_tensors": 108864,
+  "bytes": 19388363261,
+  "files": [
+    {
+      "path": "LICENSE",
+      "bytes": 1070,
+      "sha256": "30b85b6b9659f2e78aa259f8faf5d920a68dee7c9ced3fa6dba1f19f2bc4fca1"
+    },
+    {
+      "path": "README.md",
+      "bytes": 7243,
+      "sha256": "3f0894b80aefb75c2afbaa0d4fb0f0f0993f14cb9700c6a5077ea6c201637b7b"
+    },
+    {
+      "path": "amax.safetensors",
+      "bytes": 9779904,
+      "sha256": "db629a0e7cc2b77ea8a4ec560d326ae94c2af561f4fdbd39d8f611818da43d2b"
+    },
+    {
+      "path": "amax_checkpoint.json",
+      "bytes": 1137,
+      "sha256": "e2ab014bc3354ad679c4057c5890e786a45dd6517b61d049a652ac76ec9cba6c"
+    },
+    {
+      "path": "amax_checkpoint.safetensors",
+      "bytes": 9780048,
+      "sha256": "d3ef0aa55826c7f2a29c65a1d8b32368a26a527cf36f27c1665bd0bec5d34fa1"
+    },
+    {
+      "path": "chat_template.jinja",
+      "bytes": 10644,
+      "sha256": "34d5ee66b12fa6446cdae131c352b8f68cd85369e0e6fda115583805fada3891"
+    },
+    {
+      "path": "config.json",
+      "bytes": 15761,
+      "sha256": "676382abd1e90a6c85f0c8f33d45441ecd45fd514fd7b63ce5610e732d8e4996"
+    },
+    {
+      "path": "generation_config.json",
+      "bytes": 2233,
+      "sha256": "80475d7d0e56cf2729f1e0e60bfec01bda743f3829bc5c42e14f54fb6d52e612"
+    },
+    {
+      "path": "hf_quant_config.json",
+      "bytes": 8052,
+      "sha256": "9c084477c9fe496929a15dde8fb795d13b143e32911e613fe7209121b5606ee9"
+    },
+    {
+      "path": "model-hf-nonexpert-00001-of-00004.safetensors",
+      "bytes": 5356357736,
+      "sha256": "9f7ede71b2213a6962ea3801015e84edd185ddf00e1b8a0aa9ea6288c28014bf"
+    },
+    {
+      "path": "model-hf-nonexpert-00002-of-00004.safetensors",
+      "bytes": 5318341144,
+      "sha256": "b26a6aeebb246e95fcb5d54f0b6f8dd66fdfeb17fe54652ae835af1b7d64f6ef"
+    },
+    {
+      "path": "model-hf-nonexpert-00003-of-00004.safetensors",
+      "bytes": 5343414728,
+      "sha256": "0d56d129e4bd3d732968e7c5dbf3fc11e3e23566619842832325c44306cc68a5"
+    },
+    {
+      "path": "model-hf-nonexpert-00004-of-00004.safetensors",
+      "bytes": 2951938800,
+      "sha256": "43f38b4fe13a7002a5cf2841083caaac5baa65ebf72e7fb4d9c58bf9bff0280a"
+    },
+    {
+      "path": "model-inputscales.safetensors",
+      "bytes": 4726664,
+      "sha256": "4255779f031450572af8548c610fd9abfe7df89704985d18595c21788593cd05"
+    },
+    {
+      "path": "model-mtp-nonexpert.safetensors",
+      "bytes": 369674104,
+      "sha256": "31fc8804fce4cb1e0207d9fb13c19a892d85b50a7f769e763cb8d9e5b6a2e6ba"
+    },
+    {
+      "path": "model.safetensors.index.json",
+      "bytes": 4084881,
+      "sha256": "7edc92ec0cf197039a40235bd4f4906c6ecad85db7295d51479f51d859df310d"
+    },
+    {
+      "path": "processor_config.json",
+      "bytes": 909,
+      "sha256": "aae38374c94b08cc9b0547c6e64f05b951bd9735cea571c6988f5ed552bed3ed"
+    },
+    {
+      "path": "tokenizer.json",
+      "bytes": 20217442,
+      "sha256": "19e773648cb4e65de8660ea6365e10acca112d42a854923df93db4a6f333a82d"
+    },
+    {
+      "path": "tokenizer_config.json",
+      "bytes": 761,
+      "sha256": "98b1271574f41abf89427ae2dda030d94dc9478f0edc5a8bd240db213c6fd5fc"
+    }
+  ],
+  "layout": "carrier/ under checkpoint root; MODEL_ROOT points to carrier/",
+  "native_execution_tested": false,
+  "mtp_expert_tensors_replaced": 1728,
+  "mtp_nonexpert_tensors_preserved": 25,
+  "mtp_nonexpert_tensor_equality": "PASS",
+  "derivation": "MTP45 FP8 experts and scales replaced by Trellis sidecars; other tensors unchanged"
+}

evidence/prepare-mtp-inputs.py

@@ -0,0 +1,73 @@
+"""Prepare the existing BF16 MTP captures for the owner's new Trellis MTP arm."""
+import hashlib
+import json
+import os
+from pathlib import Path
+
+import numpy as np
+import torch
+from huggingface_hub import snapshot_download
+from safetensors.torch import save_file
+
+root = Path(os.environ["CV2_ROOT"])
+out = Path(os.environ["CV2_OUT"])
+dest = out / "mtp-inputs"
+dest.mkdir(exist_ok=True)
+repo = "brandonmusic/GLM-5.3-Flash-BF16-Teacher-Logits"
+rev = "95f4fdd94bf29989db2e0d1054e4931f55edb6aa"
+snapshot_download(repo, repo_type="dataset", revision=rev, local_dir=dest / "source",
+                  allow_patterns=["calibration/mtp45-ep4-full/*"], max_workers=4)
+source = dest / "source/calibration/mtp45-ep4-full"
+manifest = json.loads((source / "capture-manifest.json").read_text())
+receipt = json.loads((source / "capture-receipt.json").read_text())
+assert receipt["complete"] is True
+assert hashlib.sha256((source / "capture-manifest.json").read_bytes()).hexdigest() == receipt["capture_manifest_file_sha256"]
+assert manifest["layer"] == 45 and manifest["geometry"]["experts"] == 288
+assert manifest["model_revision"] == "a6c167b62691b2bac901344b65cb651a70f53e43"
+for info in manifest["files"].values():
+    path = source / info["path"]
+    assert path.stat().st_size == info["bytes"]
+    with path.open("rb") as stream:
+        assert hashlib.file_digest(stream, "sha256").hexdigest() == info["sha256"]
+roles = json.loads((root / "data/roles-v5.json").read_text())["roles"]["fit"]
+assert len(roles) == 64
+role_map = {r["id"]: r for r in roles}
+windows = [w for w in manifest["windows"] if w["window_id"] in role_map]
+assert len(windows) == 64
+offset = 0
+offsets = {}
+for window in manifest["windows"]:
+    offsets[window["window_id"]] = offset
+    offset += window["rows"]
+assert offset == manifest["rows"]
+for window in windows:
+    assert window["role"] == "fit" and window["rows"] == 2047
+    assert window["token_ids_sha256"] == role_map[window["window_id"]]["input_sha256"]
+streams = [("hidden_bf16", "<u2", 4096, "hidden.bf16.bin"),
+           ("topk_ids_u16le", "<u2", 8, "topk_ids.u16le.bin"),
+           ("topk_weights_f32le", "<f4", 8, "topk_weights.f32le.bin")]
+for key, dtype, width, name in streams:
+    array = np.memmap(source / manifest["files"][key]["path"], mode="r", dtype=dtype,
+                      shape=(manifest["rows"], width))
+    with (dest / name).open("wb") as stream:
+        for window in windows:
+            start = offsets[window["window_id"]]
+            stream.write(np.asarray(array[start:start + window["rows"]]).tobytes())
+# No prior Trellis/EXL3 MTP scale vectors exist in the supplied r27 overlay.
+# New signed-unit vectors provide an orthogonal coupled H512/H128 transform;
+# the unchanged codec-v2 block-scale refit handles weight magnitudes.
+generator = torch.Generator(device="cpu").manual_seed(530045)
+def signs(shape):
+    return torch.randint(0, 2, shape, generator=generator).mul_(2).sub_(1).to(torch.float16).contiguous()
+scales = {"gate_up_suh": signs((4096,)), "down_svh": signs((4096,)),
+          "gate_svh": signs((288, 2048)), "up_svh": signs((288, 2048)),
+          "down_suh": signs((288, 2048))}
+save_file(scales, str(dest / "coupled-scales.safetensors"))
+info = {"schema": "trellismx-codecv2-mtp-inputs.v1", "layer": 45,
+        "capture_repo": repo, "capture_revision": rev, "source_capture_sha256": manifest["capture_sha256"],
+        "windows": windows, "rows_per_window": 2047, "rows": 64 * 2047,
+        "scale_policy": "signed-unit coupled vectors; CPU torch seed 530045; no prior r27 MTP scales",
+        "scale_sha256": hashlib.sha256((dest / "coupled-scales.safetensors").read_bytes()).hexdigest()}
+(dest / "inputs.json").write_text(json.dumps(info, indent=2) + "\n")
+(out / "status/mtp-inputs.ready").write_text(info["scale_sha256"] + "\n")
+print(json.dumps({k: v for k, v in info.items() if k != "windows"}), flush=True)

runtime/Dockerfile

@@ -0,0 +1,6 @@
+FROM verdictai/trellismx@sha256:a0392e1c370eb933d87511928ea3aba6d5d63965cfe6eeaf9cc026a667e98ddc
+COPY patch-manifest.json install-mtp-runtime.py /opt/codecv2-mtp/
+COPY patches /opt/codecv2-mtp/patches
+RUN python3 /opt/codecv2-mtp/install-mtp-runtime.py
+ENV MODEL_ROOT=/model/carrier VLLM_TRELLISMX_CHECKPOINT=/model
+ENTRYPOINT ["/bin/bash", "/release/serve-rp2.sh"]

runtime/LICENSE-APACHE-2.0

@@ -0,0 +1,201 @@
+                                 Apache License
+                           Version 2.0, January 2004
+                        http://www.apache.org/licenses/
+
+   TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION
+
+   1. Definitions.
+
+      "License" shall mean the terms and conditions for use, reproduction,
+      and distribution as defined by Sections 1 through 9 of this document.
+
+      "Licensor" shall mean the copyright owner or entity authorized by
+      the copyright owner that is granting the License.
+
+      "Legal Entity" shall mean the union of the acting entity and all
+      other entities that control, are controlled by, or are under common
+      control with that entity. For the purposes of this definition,
+      "control" means (i) the power, direct or indirect, to cause the
+      direction or management of such entity, whether by contract or
+      otherwise, or (ii) ownership of fifty percent (50%) or more of the
+      outstanding shares, or (iii) beneficial ownership of such entity.
+
+      "You" (or "Your") shall mean an individual or Legal Entity
+      exercising permissions granted by this License.
+
+      "Source" form shall mean the preferred form for making modifications,
+      including but not limited to software source code, documentation
+      source, and configuration files.
+
+      "Object" form shall mean any form resulting from mechanical
+      transformation or translation of a Source form, including but
+      not limited to compiled object code, generated documentation,
+      and conversions to other media types.
+
+      "Work" shall mean the work of authorship, whether in Source or
+      Object form, made available under the License, as indicated by a
+      copyright notice that is included in or attached to the work
+      (an example is provided in the Appendix below).
+
+      "Derivative Works" shall mean any work, whether in Source or Object
+      form, that is based on (or derived from) the Work and for which the
+      editorial revisions, annotations, elaborations, or other modifications
+      represent, as a whole, an original work of authorship. For the purposes
+      of this License, Derivative Works shall not include works that remain
+      separable from, or merely link (or bind by name) to the interfaces of,
+      the Work and Derivative Works thereof.
+
+      "Contribution" shall mean any work of authorship, including
+      the original version of the Work and any modifications or additions
+      to that Work or Derivative Works thereof, that is intentionally
+      submitted to Licensor for inclusion in the Work by the copyright owner
+      or by an individual or Legal Entity authorized to submit on behalf of
+      the copyright owner. For the purposes of this definition, "submitted"
+      means any form of electronic, verbal, or written communication sent
+      to the Licensor or its representatives, including but not limited to
+      communication on electronic mailing lists, source code control systems,
+      and issue tracking systems that are managed by, or on behalf of, the
+      Licensor for the purpose of discussing and improving the Work, but
+      excluding communication that is conspicuously marked or otherwise
+      designated in writing by the copyright owner as "Not a Contribution."
+
+      "Contributor" shall mean Licensor and any individual or Legal Entity
+      on behalf of whom a Contribution has been received by Licensor and
+      subsequently incorporated within the Work.
+
+   2. Grant of Copyright License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      copyright license to reproduce, prepare Derivative Works of,
+      publicly display, publicly perform, sublicense, and distribute the
+      Work and such Derivative Works in Source or Object form.
+
+   3. Grant of Patent License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      (except as stated in this section) patent license to make, have made,
+      use, offer to sell, sell, import, and otherwise transfer the Work,
+      where such license applies only to those patent claims licensable
+      by such Contributor that are necessarily infringed by their
+      Contribution(s) alone or by combination of their Contribution(s)
+      with the Work to which such Contribution(s) was submitted. If You
+      institute patent litigation against any entity (including a
+      cross-claim or counterclaim in a lawsuit) alleging that the Work
+      or a Contribution incorporated within the Work constitutes direct
+      or contributory patent infringement, then any patent licenses
+      granted to You under this License for that Work shall terminate
+      as of the date such litigation is filed.
+
+   4. Redistribution. You may reproduce and distribute copies of the
+      Work or Derivative Works thereof in any medium, with or without
+      modifications, and in Source or Object form, provided that You
+      meet the following conditions:
+
+      (a) You must give any other recipients of the Work or
+          Derivative Works a copy of this License; and
+
+      (b) You must cause any modified files to carry prominent notices
+          stating that You changed the files; and
+
+      (c) You must retain, in the Source form of any Derivative Works
+          that You distribute, all copyright, patent, trademark, and
+          attribution notices from the Source form of the Work,
+          excluding those notices that do not pertain to any part of
+          the Derivative Works; and
+
+      (d) If the Work includes a "NOTICE" text file as part of its
+          distribution, then any Derivative Works that You distribute must
+          include a readable copy of the attribution notices contained
+          within such NOTICE file, excluding those notices that do not
+          pertain to any part of the Derivative Works, in at least one
+          of the following places: within a NOTICE text file distributed
+          as part of the Derivative Works; within the Source form or
+          documentation, if provided along with the Derivative Works; or,
+          within a display generated by the Derivative Works, if and
+          wherever such third-party notices normally appear. The contents
+          of the NOTICE file are for informational purposes only and
+          do not modify the License. You may add Your own attribution
+          notices within Derivative Works that You distribute, alongside
+          or as an addendum to the NOTICE text from the Work, provided
+          that such additional attribution notices cannot be construed
+          as modifying the License.
+
+      You may add Your own copyright statement to Your modifications and
+      may provide additional or different license terms and conditions
+      for use, reproduction, or distribution of Your modifications, or
+      for any such Derivative Works as a whole, provided Your use,
+      reproduction, and distribution of the Work otherwise complies with
+      the conditions stated in this License.
+
+   5. Submission of Contributions. Unless You explicitly state otherwise,
+      any Contribution intentionally submitted for inclusion in the Work
+      by You to the Licensor shall be under the terms and conditions of
+      this License, without any additional terms or conditions.
+      Notwithstanding the above, nothing herein shall supersede or modify
+      the terms of any separate license agreement you may have executed
+      with Licensor regarding such Contributions.
+
+   6. Trademarks. This License does not grant permission to use the trade
+      names, trademarks, service marks, or product names of the Licensor,
+      except as required for reasonable and customary use in describing the
+      origin of the Work and reproducing the content of the NOTICE file.
+
+   7. Disclaimer of Warranty. Unless required by applicable law or
+      agreed to in writing, Licensor provides the Work (and each
+      Contributor provides its Contributions) on an "AS IS" BASIS,
+      WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or
+      implied, including, without limitation, any warranties or conditions
+      of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A
+      PARTICULAR PURPOSE. You are solely responsible for determining the
+      appropriateness of using or redistributing the Work and assume any
+      risks associated with Your exercise of permissions under this License.
+
+   8. Limitation of Liability. In no event and under no legal theory,
+      whether in tort (including negligence), contract, or otherwise,
+      unless required by applicable law (such as deliberate and grossly
+      negligent acts) or agreed to in writing, shall any Contributor be
+      liable to You for damages, including any direct, indirect, special,
+      incidental, or consequential damages of any character arising as a
+      result of this License or out of the use or inability to use the
+      Work (including but not limited to damages for loss of goodwill,
+      work stoppage, computer failure or malfunction, or any and all
+      other commercial damages or losses), even if such Contributor
+      has been advised of the possibility of such damages.
+
+   9. Accepting Warranty or Additional Liability. While redistributing
+      the Work or Derivative Works thereof, You may choose to offer,
+      and charge a fee for, acceptance of support, warranty, indemnity,
+      or other liability obligations and/or rights consistent with this
+      License. However, in accepting such obligations, You may act only
+      on Your own behalf and on Your sole responsibility, not on behalf
+      of any other Contributor, and only if You agree to indemnify,
+      defend, and hold each Contributor harmless for any liability
+      incurred by, or claims asserted against, such Contributor by reason
+      of your accepting any such warranty or additional liability.
+
+   END OF TERMS AND CONDITIONS
+
+   APPENDIX: How to apply the Apache License to your work.
+
+      To apply the Apache License to your work, attach the following
+      boilerplate notice, with the fields enclosed by brackets "[]"
+      replaced with your own identifying information. (Don't include
+      the brackets!)  The text should be enclosed in the appropriate
+      comment syntax for the file format. We also recommend that a
+      file or class name and description of purpose be included on the
+      same "printed page" as the copyright notice for easier
+      identification within third-party archives.
+
+   Copyright [yyyy] [name of copyright owner]
+
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+
+       http://www.apache.org/licenses/LICENSE-2.0
+
+   Unless required by applicable law or agreed to in writing, software
+   distributed under the License is distributed on an "AS IS" BASIS,
+   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+   See the License for the specific language governing permissions and
+   limitations under the License.

runtime/install-mtp-runtime.py

@@ -0,0 +1,26 @@
+"""Install the three pinned MTP loader changes inside the existing RP2 image."""
+import hashlib
+import json
+import os
+from pathlib import Path
+
+root = Path(__file__).resolve().parent
+manifest = json.loads((root / "patch-manifest.json").read_text())
+pending = []
+for record in manifest["files"]:
+    target = Path(record["target"])
+    data = (root / "patches" / record["source"]).read_bytes()
+    assert hashlib.sha256(data).hexdigest() == record["patched_sha256"]
+    actual = hashlib.sha256(target.read_bytes()).hexdigest()
+    if actual == record["patched_sha256"]:
+        continue
+    if actual != record["original_sha256"]:
+        raise RuntimeError(f"Unsupported runtime source at {target}; use {manifest['base_image']}")
+    compile(data, str(target), "exec")
+    pending.append((target, data))
+for target, data in pending:
+    temporary = target.with_suffix(".codecv2-mtp.tmp")
+    temporary.write_bytes(data)
+    temporary.chmod(target.stat().st_mode & 0o777)
+    os.replace(temporary, target)
+print("TrellisMX MTP45 runtime installed; existing main-model kernels preserved", flush=True)

runtime/patch-manifest.json

@@ -0,0 +1,23 @@
+{
+  "base_image": "verdictai/trellismx@sha256:a0392e1c370eb933d87511928ea3aba6d5d63965cfe6eeaf9cc026a667e98ddc",
+  "files": [
+    {
+      "source": "vllm_utils_trellismx.py",
+      "target": "/opt/glm53-flash/vllm/vllm/utils/trellismx.py",
+      "original_sha256": "27be1af286b3fe6eee1fa11602f9fd87142ffa0e13bb9895a9a96346beb91fa3",
+      "patched_sha256": "c228f6c52c8f34cb63d37809cea03a944c1fcdbf013459434f29c95eb43d821f"
+    },
+    {
+      "source": "vllm_quant_trellismx.py",
+      "target": "/opt/glm53-flash/vllm/vllm/model_executor/layers/quantization/trellismx.py",
+      "original_sha256": "26512cae767be9b0f03a0dfbe2280b8f8778b7889eca146e03fbc6e740a3568e",
+      "patched_sha256": "6d849966ed3ea4b5f32d08702d58d20f6b24985de5a65f100ef76ee679ea267c"
+    },
+    {
+      "source": "p8_native_kernel.py",
+      "target": "/opt/glm53-flash/b12x/b12x/moe/_shared/trellismx/p8_native_kernel.py",
+      "original_sha256": "fc88f5ed2a466bbca8459b83dadfd14d5063a4e19f45626a31ca6f7c52c02078",
+      "patched_sha256": "6f3bb3c2c61f1881b2e0420a6606fb0cc3d8c37b7939ad64c15a043bfdb18167"
+    }
+  ]
+}

runtime/patches/p8_native_kernel.py

@@ -0,0 +1,887 @@
+"""Experimental TP-local P8 procedural-MCG MoE runtime for GLM-5.3.
+
+This is a direct device path: a K3, K4 or K5 trellis stream is decoded to E4M3 inside
+the MMA kernel and physical UE8M0/32 scales are consumed by the tensor core.
+It implements the frozen TP4 identity-boundary P8 contract and an opt-in M1
+H128 suh/svh scale component for a GLM routed layer whose sidecar carries the
+matching immutable layer identity.
+"""
+from __future__ import annotations
+
+from dataclasses import dataclass
+from pathlib import Path
+import os
+
+import cutlass
+import cutlass.cute as cute
+import torch
+from cutlass.base_dsl.compiler import OptLevel
+from cutlass.cute.runtime import make_ptr
+from safetensors import safe_open
+from .p8_coupled_scales import (
+    COUPLED_SCHEMA as P8_COUPLED_SCHEMA,
+    SCALE_NAMES,
+    SCHEMA as P8_SCALE_COMPONENT_SCHEMA,
+    validate_coupled_component,
+    validate_scale_component,
+)
+from .policy_smallm_schedule import P8SmallMGeometry, p8_small_m_scratch_layout, use_small_m
+from .tile_policy import select_tile
+
+from b12x._lib.compiler import KernelCompileSpec, compile as b12x_compile
+from b12x._lib.utils import get_max_active_clusters
+from b12x.moe._shared.kernels.route_hoist_dynamic import MoEDynamicKernelBackend
+from b12x.moe.fused_moe._impl import (
+    _DynamicMoEW4A8Launch,
+    _e8m0_scale_to_w4a8_sfb_inplace,
+    _launch_dynamic_topk_sum,
+    current_cuda_stream,
+)
+
+
+def _gptr(dtype, tensor: torch.Tensor, align: int = 16):
+    return make_ptr(
+        dtype, tensor.data_ptr(), cute.AddressSpace.gmem, assumed_align=align
+    )
+
+
+def _fake_i32(shape: tuple[int, ...]):
+    return cute.runtime.make_fake_compact_tensor(
+        cutlass.Int32, shape, assumed_align=4
+    )
+
+
+def _fake_f32(shape: tuple[int, ...]):
+    return cute.runtime.make_fake_compact_tensor(
+        cutlass.Float32, shape, assumed_align=16
+    )
+
+
+@dataclass
+class _CompiledArm:
+    compiled: object
+    tile_m: int
+    materialized: bool
+    mac: int
+
+
+class P8NativeTPMoE:
+    """Own one TP rank's physical P8 payload and launch compiled MoE kernels."""
+
+    def __init__(
+        self,
+        sidecar: Path | tuple[Path, Path],
+        *,
+        device: torch.device,
+        tp_rank: int,
+        layer: int = 3,
+        expected_design_sha256: str | None = None,
+        expected_transform_sha256: str | None = None,
+        topk: int = 8,
+        hidden: int = 4096,
+        intermediate: int = 512,
+        swiglu_limit: float = 10.0,
+        force_materialized: bool | None = None,
+        mac_override: int | None = None,
+        deterministic_output: bool = True,
+        small_m_scheduler: bool = False,
+        fc1_tile_n: int = 128,
+        debug_capture: bool = False,
+        diagnostic_raw_fc1: bool = False,
+        fuse_scratch_zero: bool = False,
+        compact_scale_storage: bool = False,
+        compact_input_storage: bool = False,
+        shared_workspace: bool = False,
+        world_size: int = 4,
+        tp4_parent_sha256: tuple[str, str] | None = None,
+        prefill_chunk_tokens: int = 0,
+        grid_policy: bool | None = None,
+        grouped_m16: bool = False,
+        fuse_grouped_scratch: bool = False,
+        tile_major_tasks: bool = False,
+        fc1_pipeline_stages: int = 2,
+        fc1_warps: int = 4,
+        fc1_a_swizzle_rotate: bool = False,
+        fc1_warp_quant: bool | None = None,
+        fc1_exact_staging: bool = False,
+        fc1_broadcast_a: bool | None = None,
+    ) -> None:
+        self.device = torch.device(device)
+        # D-x2 (PREREG amendment 2): captured once; joins the compile spec below.
+        self.p8_down_remainder = os.environ.get("B12X_P8_DOWN_REMAINDER") == "1"
+        self._dx2_planes = 2 if self.p8_down_remainder else 1
+        self._dx2_phases = os.environ.get("B12X_P8_DOWN_REMAINDER_PHASES", "both")
+        self.p8_dx2_rowpack = os.environ.get("B12X_P8_DOWN_REMAINDER") in ("rp", "rp2")
+        self.p8_input_rowpack = os.environ.get("B12X_P8_DOWN_REMAINDER") == "rp2"
+        # EPI-PAR: bit-identical parallel small-M FC1 epilogue; joins the compile spec below.
+        self.p8_epi_par = os.environ.get("B12X_P8_EPI_PAR") == "1"
+        if self.p8_epi_par:
+            print(f"P8_EPI_PAR_ACTIVE layer={layer} rank={tp_rank}", flush=True)
+        if self.p8_dx2_rowpack:
+            print(f"P8_DX2_ROWPACK_ACTIVE layer={layer} rank={tp_rank} input_hop={int(self.p8_input_rowpack)}", flush=True)
+        self.grouped_m16 = bool(grouped_m16)
+        self.fuse_grouped_scratch = bool(fuse_grouped_scratch)
+        self.tile_major_tasks = bool(tile_major_tasks)
+        if fc1_pipeline_stages not in (2, 3):
+            raise ValueError('FC1 pipeline supports two or three stages')
+        self.fc1_pipeline_stages = fc1_pipeline_stages
+        if fc1_warps not in (4, 8):
+            raise ValueError('FC1 supports four or eight warps')
+        self.fc1_warps = fc1_warps
+        self.fc1_exact_staging = bool(fc1_exact_staging)
+        self.fc1_broadcast_a = (os.environ.get('GLM53_P8_FC1_BROADCAST_A') == '1'
+                                if fc1_broadcast_a is None else bool(fc1_broadcast_a))
+        self.fc1_a_swizzle_rotate = bool(fc1_a_swizzle_rotate)
+        self.fc1_warp_quant = (os.environ.get('GLM53_P8_FC1_WARP_QUANT') == '1'
+                              if fc1_warp_quant is None else bool(fc1_warp_quant))
+        self.grid_policy = (os.environ.get('GLM53_P8_GRID_POLICY') == '1'
+                            if grid_policy is None else bool(grid_policy))
+        self.prefill_chunk_tokens = int(prefill_chunk_tokens)
+        if self.prefill_chunk_tokens < 0 or self.prefill_chunk_tokens % 64:
+            raise ValueError("P8 prefill chunk must be zero or a positive multiple of 64")
+        self.compact_scale_storage = bool(compact_scale_storage)
+        self.compact_input_storage = bool(compact_input_storage)
+        self.shared_workspace = bool(shared_workspace)
+        self.tp_rank = int(tp_rank)
+        self.world_size = int(world_size)
+        if self.world_size != 4 or self.tp_rank not in range(4):
+            raise ValueError("P8 native supports TP4 only; TP2 validators are unsupported")
+        self.layer = int(layer)
+        if not 3 <= self.layer <= 45:
+            raise ValueError("P8 native layer must be in GLM routed layers 3..44 or MTP45")
+        self.topk = int(topk)
+        self.hidden = int(hidden)
+        self.intermediate = int(intermediate)
+        self.swiglu_limit = float(swiglu_limit)
+        self.force_materialized = force_materialized
+        self.mac_override = None if mac_override is None else int(mac_override)
+        self.deterministic_output = bool(deterministic_output)
+        # Explicit developmental opt-in. M2/M3 and prefill retain baseline
+        # selection; this is not enabled through a serving environment flag.
+        self.small_m_scheduler = bool(small_m_scheduler)
+        self.fc1_tile_n = int(fc1_tile_n)
+        self.debug_capture = bool(debug_capture)
+        self.diagnostic_raw_fc1 = bool(diagnostic_raw_fc1)
+        self.fuse_scratch_zero = bool(fuse_scratch_zero)
+        self._scratch_layout = (
+            p8_small_m_scratch_layout(intermediate=self.intermediate) if self.fuse_scratch_zero else None
+        )
+        self.debug_tensors = {}
+        if self.fc1_tile_n not in (32, 64, 128):
+            raise ValueError("FC1 tile N must be 32, 64, or 128")
+        if self.fc1_tile_n != 128 and (
+            not self.small_m_scheduler or self.swiglu_limit != 10.0
+        ):
+            raise ValueError("Narrow FC1 requires small-M and SwiGLU limit 10")
+        if self.small_m_scheduler and (
+            not self.deterministic_output or force_materialized is not None
+            or (topk, hidden, intermediate) != (8, 4096, 2048 // self.world_size)
+        ):
+            raise ValueError("P8 small-M requires deterministic GLM TP4 and automatic fallback")
+        if self.diagnostic_raw_fc1 and (
+            not self.debug_capture
+            or not self.small_m_scheduler
+            or self.fc1_tile_n != 128
+            or self.fuse_scratch_zero
+        ):
+            raise ValueError(
+                "raw FC1 diagnostic requires debug M1 small-M N128 capture"
+            )
+        if self.mac_override is not None and self.mac_override <= 0:
+            raise ValueError("mac_override must be positive")
+        if tp4_parent_sha256 is not None:
+            if self.world_size != 2:
+                raise ValueError("Parent-pair adapter requires TP2")
+            from glm53_nvfp4.p8_tp2_repack import open_tp2_pair
+            source = open_tp2_pair(sidecar, tp4_parent_sha256, layer=self.layer, rank=self.tp_rank)
+        else:
+            source = safe_open(sidecar, framework="pt", device="cpu")
+        with source as src:
+            metadata = src.metadata() or {}
+            schema = metadata.get("schema")
+            source_design_sha256 = metadata.get("source_design_sha256")
+            bits_text = metadata.get("bits", "")
+            if bits_text not in {"3", "4", "5"}:
+                raise RuntimeError(f"invalid P8 trellis rate: {bits_text!r}")
+            self.trellis_bits = int(bits_text)
+            base_required = {
+                "layer": str(self.layer),
+                "rank": str(self.tp_rank),
+                "world_size": str(self.world_size),
+                "bits": bits_text,
+                "alphabet": "e4m3",
+                "scale": "ue8m0-k32",
+                "law": "procedural-mcg-alpha2",
+                "ldlq": "false",
+            }
+            identity_schema = schema in {
+                    "glm53-p8-identity-mcg-tp4-rank.v1",
+                    "glm53-p8-mcg-tp4-rank.v2",
+            }
+            scale_component_schema = schema == P8_SCALE_COMPONENT_SCHEMA
+            full_coupled_schema = schema == P8_COUPLED_SCHEMA.replace("tp4", f"tp{self.world_size}")
+            if self.world_size == 2 and not full_coupled_schema:
+                raise RuntimeError("TP2 port only supports the full-coupled schema")
+            if (
+                not (identity_schema or scale_component_schema or full_coupled_schema)
+                or any(metadata.get(key) != value for key, value in base_required.items())
+                or (identity_schema and metadata.get("boundary") != "identity")
+            ):
+                raise RuntimeError(f"invalid P8 native sidecar metadata: {metadata}")
+            if full_coupled_schema or schema in {
+                "glm53-p8-mcg-tp4-rank.v2",
+                P8_SCALE_COMPONENT_SCHEMA,
+                P8_COUPLED_SCHEMA,
+            }:
+                if (
+                    not isinstance(source_design_sha256, str)
+                    or len(source_design_sha256) != 64
+                    or any(char not in "0123456789abcdef" for char in source_design_sha256)
+                ):
+                    raise RuntimeError("v2 P8 sidecar lacks a valid source design hash")
+                if (
+                    expected_design_sha256 is not None
+                    and source_design_sha256 != expected_design_sha256
+                ):
+                    raise RuntimeError("P8 sidecar does not match the expected design")
+            elif expected_design_sha256 is not None:
+                raise RuntimeError("historical P8 sidecars cannot satisfy a v2 design pin")
+            w13 = src.get_tensor("w13_trellis")
+            w2 = src.get_tensor("w2_trellis")
+            w13_scale = src.get_tensor("w13_scale_ue8m0")
+            w2_scale = src.get_tensor("w2_scale_ue8m0")
+            scale_tensors = (
+                {name: src.get_tensor(name) for name in SCALE_NAMES}
+                if scale_component_schema or full_coupled_schema
+                else None
+            )
+        self.source_design_sha256 = source_design_sha256
+        experts = int(w13.shape[1])
+        stream_words = 16 * self.trellis_bits
+        if tuple(w13.shape) != (
+            2, experts, hidden // 16, intermediate // 16, stream_words
+        ):
+            raise RuntimeError(f"unexpected W13 trellis shape {tuple(w13.shape)}")
+        if tuple(w2.shape) != (
+            experts, intermediate // 16, hidden // 16, stream_words
+        ):
+            raise RuntimeError(f"unexpected W2 trellis shape {tuple(w2.shape)}")
+        if tuple(w13_scale.shape) != (experts, 2 * intermediate, hidden // 32):
+            raise RuntimeError(f"unexpected W13 scale shape {tuple(w13_scale.shape)}")
+        if tuple(w2_scale.shape) != (experts, hidden, intermediate // 32):
+            raise RuntimeError(f"unexpected W2 scale shape {tuple(w2_scale.shape)}")
+        self.experts = experts
+        self.scale_component = None
+        self.full_coupled = bool(full_coupled_schema)
+        if (self.grouped_m16 or self.fuse_grouped_scratch) and not self.full_coupled:
+            raise ValueError("grouped scratch requires the full-coupled kernel owner")
+        if self.fuse_grouped_scratch and not self.grouped_m16:
+            raise ValueError("fused grouped scratch requires grouped_m16")
+        if self.shared_workspace and (not self.full_coupled or self.debug_capture or self.fuse_scratch_zero):
+            raise ValueError('shared workspace requires full coupling without retained debug tensors or fused arena')
+        if self.compact_scale_storage and not self.full_coupled:
+            raise ValueError("compact scales require the external full-coupled owners")
+        if self.full_coupled and expected_transform_sha256 is None:
+            raise RuntimeError(
+                "full-coupled P8 requires an externally pinned encoder transform"
+            )
+        if scale_tensors is not None:
+            validator = (
+                validate_coupled_component
+                if self.full_coupled
+                else validate_scale_component
+            )
+            # Preserve compatibility with the pinned TP4 validator during
+            # isolated storage diagnostics. TP2 requires the ported validator.
+            validator_kwargs = {} if self.world_size == 4 else {"world_size": self.world_size}
+            if self.full_coupled:
+                validator_kwargs["expected_transform_sha256"] = (
+                    expected_transform_sha256
+                )
+            self.scale_component = validator(
+                metadata, scale_tensors, layer=self.layer, rank=self.tp_rank,
+                experts=experts, hidden=hidden, intermediate=intermediate,
+                **validator_kwargs,
+            )
+            if not self.small_m_scheduler or self.fc1_tile_n != 128:
+                raise RuntimeError(
+                    "P8 scale component requires the M1 N128 owner path"
+                )
+        # The trellis storage is byte-for-byte the same size as the packed
+        # E2M1 descriptor carrier expected by the inherited W4A8 launch ABI.
+        # Alias it for the descriptor-only arguments instead of allocating a
+        # second ~0.9 GiB of unread dummy weights per layer and TP rank.  The
+        # kernel reads the procedural stream through the uint32 pointers below;
+        # it never dereferences the descriptor carrier values.
+        w13_stream_storage = w13.to(device=self.device).contiguous()
+        w2_stream_storage = w2.to(device=self.device).contiguous()
+        self.w13_stream = w13_stream_storage.view(torch.int32).reshape(-1)
+        if self.fc1_exact_staging:
+            if ((self.world_size, self.experts, self.hidden, self.intermediate, self.fc1_tile_n)
+                    != (4, 288, 4096, 512, 128) or not self.full_coupled):
+                raise ValueError('exact staging requires full-coupled TP4 E288/H4096/I512 N128')
+            expected_words = 2 * 288 * 4096 * 512 * self.trellis_bits // 32
+            if self.w13_stream.numel() != expected_words:
+                raise ValueError('exact staging FC1 stream extent mismatch')
+        self.w2_stream = w2_stream_storage.view(torch.int32).reshape(-1)
+        w13_scale = w13_scale.to(device=self.device).contiguous()
+        w2_scale = w2_scale.to(device=self.device).contiguous()
+        # The monolithic kernel consumes the logical [E, N, K/32] UE8M0
+        # plane through the sfb_*_mx ABI slots.  The split materialized
+        # kernels consume a separately repacked copy through *_sfb_rp.
+        # Keep both representations: a one-byte sentinel in the logical slots
+        # is an out-of-bounds scale read, not an identity scale.
+        self.w13_scale_mx = w13_scale.reshape(-1)
+        self.w2_scale_mx = w2_scale.reshape(-1)
+        self.w13_sfb = _e8m0_scale_to_w4a8_sfb_inplace(
+            w13_scale.clone(),
+            weight_E=experts,
+            rows=2 * intermediate,
+            k_dim=hidden,
+            gated_half_rows=intermediate,
+        ).reshape(-1)
+        self.w2_sfb = _e8m0_scale_to_w4a8_sfb_inplace(
+            w2_scale.clone(),
+            weight_E=experts,
+            rows=hidden,
+            k_dim=intermediate,
+        ).reshape(-1)
+        if self.compact_scale_storage:
+            # Full coupling forces external materialized FC1/FC2 for every M.
+            # Logical-scale ABI arguments remain non-null aliases, but are
+            # not read by those owners. Never enable for a monolithic arm.
+            # This is opt-in pending device closure against separate storage.
+            self.w13_scale_mx = self.w13_sfb
+            self.w2_scale_mx = self.w2_sfb
+        # These are descriptor carriers only; they alias the trellis storage
+        # above and therefore add zero payload bytes.  Their trailing extent is
+        # the PACKED row length, which is bits/8 bytes per weight: hidden // 2
+        # only at K4.  Deriving it from the stored rate keeps K4 byte-identical
+        # while letting K3 and K5 describe their own shorter or longer rows.
+        w13_row_bytes = hidden * self.trellis_bits // 8
+        w2_row_bytes = intermediate * self.trellis_bits // 8
+        w13_dummy_bytes = experts * 2 * intermediate * w13_row_bytes
+        w2_dummy_bytes = experts * hidden * w2_row_bytes
+        self.w13_dummy = w13_stream_storage.view(torch.uint8).reshape(-1)[
+            :w13_dummy_bytes
+        ].reshape(
+            experts, 2 * intermediate, w13_row_bytes
+        )
+        self.w2_dummy = w2_stream_storage.view(torch.uint8).reshape(-1)[
+            :w2_dummy_bytes
+        ].reshape(
+            experts, hidden, w2_row_bytes
+        )
+        self.sentinel = torch.zeros(1, dtype=torch.uint8, device=self.device)
+        self.zero_lut = torch.zeros(1, dtype=torch.uint8, device=self.device)
+        # MCG never dereferences the LUT pointer. The diagnostic-only arm
+        # reuses that dead ABI slot for exactly 128 FP32 trace values (512 B),
+        # initialized to an all-ones NaN sentinel so partial writes fail closed.
+        self.input_prequant_trace = (
+            torch.full((512,), 0xFF, dtype=torch.uint8, device=self.device)
+            if self.diagnostic_raw_fc1
+            else self.zero_lut
+        )
+        self.zero_rotation = torch.zeros(1, dtype=torch.float16, device=self.device)
+        self.scale_component_packed = (
+            self.scale_component.packed.to(device=self.device)
+            if self.scale_component is not None
+            else self.zero_rotation
+        )
+        self.ones = torch.ones(experts, dtype=torch.float32, device=self.device)
+        if self.small_m_scheduler and self.experts != 288:
+            raise ValueError("P8 small-M requires 288 experts")
+        # v11: the small-M owner path is compiled per stored rate (K3/K4/K5);
+        # M>1 on a non-K4 layer is served row by row through that same exact
+        # kernel because the grouped M64 prefill kernels remain K4-only.
+        self._compiled: dict[tuple[bool, bool], _CompiledArm] = {}
+        self._coupled_reducer = None
+
+    def _compile(self, materialized: bool, small_m: bool = False, expected_m: int | None = None) -> _CompiledArm:
+        if self.compact_scale_storage and not (self.full_coupled and materialized):
+            raise RuntimeError("compact scales cannot enter a monolithic path")
+        selected_m = select_tile(expected_m if expected_m is not None else (1 if small_m else 4096))[0]
+        cache_key = (materialized, small_m, selected_m)
+        cached = self._compiled.get(cache_key)
+        if cached is not None:
+            return cached
+        tile_m = selected_m
+        if self.grouped_m16:
+            raise RuntimeError("fixed M16 override conflicts with requested tile policy")
+        mac = (
+            self.mac_override
+            if self.mac_override is not None
+            else (64 if materialized else int(get_max_active_clusters(1)))
+        )
+        kernel = MoEDynamicKernelBackend(
+            16,
+            (tile_m, 128),
+            activation="silu",
+            quant_recipe="w4a8_trellis",
+            w4a8_repacked=True,
+            num_topk=self.topk,
+            trellis_bits=self.trellis_bits,
+            trellis_codebook="mcg",
+            trellis_scaled=True,
+            trellis_identity_boundary=not self.full_coupled,
+            direct_routing=small_m,
+            materialize_intermediate=materialized,
+            p8_small_m=small_m,
+            p8_fc1_tile_n=self.fc1_tile_n if small_m else 128,
+            p8_scale_sandwich=self.scale_component is not None,
+            p8_full_coupled=self.full_coupled,
+            share_input_across_experts=materialized,
+            deterministic_output=self.deterministic_output,
+            swiglu_limit=self.swiglu_limit,
+        )
+        if self.p8_down_remainder:
+            if not self.full_coupled:
+                raise RuntimeError("D-x2 is implemented for the full-coupled P8 paths only")
+            if self.fc1_warp_quant:
+                raise RuntimeError("D-x2 needs fc1_warp_quant=False: the warp-quant FC1 epilogue has no remainder writer")
+            from b12x.moe._shared.kernels.p8_down_remainder import enable_down_remainder
+            enable_down_remainder(kernel, self._dx2_phases)
+        if self.p8_dx2_rowpack and small_m:
+            if not self.full_coupled:
+                raise RuntimeError("D-x2-RP is implemented for the full-coupled P8 paths only")
+            if self.fc1_warp_quant:
+                raise RuntimeError("D-x2-RP/RP2 need fc1_warp_quant=False: the warp-quant FC1 epilogue has no remainder writer")
+            if self.tile_major_tasks or self.grouped_m16:
+                raise RuntimeError("D-x2-RP/RP2 need one route per direct tile (tile_major_tasks=False, grouped_m16=False)")
+            if self.p8_input_rowpack and not self.fc1_broadcast_a:
+                raise RuntimeError("D-x2-RP2 needs fc1_broadcast_a=True: otherwise ordinary FC1 staging also writes A row 8")
+            from b12x.moe._shared.kernels.p8_down_remainder import enable_rowpack
+            enable_rowpack(kernel, input_hop=self.p8_input_rowpack)
+        if self.diagnostic_raw_fc1:
+            if not (small_m and self.full_coupled):
+                raise RuntimeError(
+                    "raw FC1 diagnostic dispatched outside full-coupled M1"
+                )
+            from b12x.moe._shared.kernels.p8_h128_fc1 import (
+                P8H128FC1RawCaptureKernel,
+            )
+
+            kernel.materialized_phase1_kernel = P8H128FC1RawCaptureKernel()
+            kernel.p8_input_prequant_diagnostic = True
+        if self.full_coupled:
+            # Both M1 and grouped owners share the scale/sign plane geometry.
+            kernel.materialized_phase1_kernel.p8_intermediate = self.intermediate
+            kernel.materialized_phase2_kernel.p8_intermediate = self.intermediate
+            if small_m:
+                kernel.materialized_phase1_kernel.p8_tile_major = self.tile_major_tasks
+                kernel.materialized_phase2_kernel.p8_tile_major = self.tile_major_tasks
+                fc1 = kernel.materialized_phase1_kernel
+                fc1.p8_a_swizzle_rotate = self.fc1_a_swizzle_rotate
+                fc1.p8_warp_quant = self.fc1_warp_quant
+                fc1.p8_epi_par = self.p8_epi_par
+                fc1.p8_exact_staging = self.fc1_exact_staging
+                # Direct owner invokes _run_task with valid_rows=1. Never
+                # apply this specialization to grouped multi-row owners.
+                fc1.p8_broadcast_a = self.fc1_broadcast_a
+                fc1.num_warps = self.fc1_warps
+                fc1.threads_per_cta = 32 * self.fc1_warps
+                fc1.p8_n8_per_warp = 16 // self.fc1_warps
+                fc1.owned_row_groups = fc1.tile_m // self.fc1_warps
+                fc1.p8_pipeline_stages = self.fc1_pipeline_stages
+                fc1.shared_bytes = max(fc1.shared_bytes, self.fc1_pipeline_stages * fc1.stage_bytes)
+                fc1.shared_words = (fc1.shared_bytes + 3) // 4
+                fc1.trellis_lut_offset = fc1.shared_bytes
+        launch = _DynamicMoEW4A8Launch(
+            kernel,
+            k=self.hidden,
+            n=self.intermediate,
+            w1_n=2 * self.intermediate,
+            num_topk=self.topk,
+        )
+
+        def ptr(dtype, address: int, align: int = 16):
+            return make_ptr(dtype, address, cute.AddressSpace.gmem, assumed_align=align)
+
+        def fake_ptr_u8():
+            return ptr(cutlass.Uint8, 16)
+
+        def fake_ptr_i32():
+            return ptr(cutlass.Int32, 4, 4)
+
+        def fake_ptr_u32():
+            return ptr(cutlass.Uint32, 16)
+
+        b_w13_fake = cute.runtime.make_fake_compact_tensor(
+            cutlass.Float4E2M1FN,
+            (2 * self.intermediate, self.hidden, self.experts),
+            stride_order=(1, 0, 2),
+            assumed_align=16,
+        )
+        b_w2_fake = cute.runtime.make_fake_compact_tensor(
+            cutlass.Float4E2M1FN,
+            (self.hidden, self.intermediate, self.experts),
+            stride_order=(1, 0, 2),
+            assumed_align=16,
+        )
+        compiled = b12x_compile(
+            launch,
+            ptr(cutlass.BFloat16, 16),
+            fake_ptr_i32(),
+            ptr(cutlass.Float32, 4, 4),
+            ptr(cutlass.Float4E2M1FN, 16),
+            ptr(cutlass.Float8E4M3FN, 16),
+            fake_ptr_u8(),
+            fake_ptr_u8(),
+            fake_ptr_u32(),
+            _fake_i32((1,)), _fake_i32((1,)), _fake_i32((1,)),
+            _fake_i32((1,)), _fake_i32((1,)), _fake_i32((1,)), _fake_i32((1,)),
+            fake_ptr_i32(), fake_ptr_i32(), fake_ptr_i32(),
+            fake_ptr_i32(), fake_ptr_i32(), fake_ptr_i32(), fake_ptr_i32(),
+            b_w13_fake,
+            ptr(cutlass.Float8E4M3FN, 16),
+            b_w2_fake,
+            ptr(cutlass.Float8E4M3FN, 16),
+            fake_ptr_u8(), fake_ptr_u8(), fake_ptr_u8(), fake_ptr_u8(),
+            fake_ptr_u32(), fake_ptr_u32(), fake_ptr_u32(), fake_ptr_u32(),
+            _fake_i32((self.experts,)),
+            _fake_i32((self.experts,)),
+            _fake_i32((self.experts + 1,)),
+            _fake_f32((self.experts,)), _fake_f32((self.experts,)),
+            _fake_f32((self.experts,)), _fake_f32((self.experts,)),
+            ptr(
+                cutlass.Float32 if self.full_coupled else cutlass.BFloat16,
+                16,
+            ),
+            fake_ptr_i32(),
+            ptr(cutlass.Float32, 16),
+            1, 1, 1, 1, 1, 1, 1,
+            current_cuda_stream(),
+            fake_ptr_u8(),
+            ptr(cutlass.Float16, 16),
+            # The compile spec is the JIT cache key and includes every specialized
+            # field, especially stored rate, topology and epilogue dimensions.
+            compile_spec=KernelCompileSpec.from_fields(
+                "glm53.p8.native.tp",
+                4,
+                ("fc1_row_alias376", 1),
+                ("requested_m_regime_direct_policy", 1),
+                ("fc1_route_hoist", 1),
+                ("tile_m", tile_m),
+                ("trellis_bits", self.trellis_bits),
+                ("mcg_k5_funnel", int(small_m and self.trellis_bits == 5)),
+                ("fc2_carveout100_grid564", int(small_m)),
+                ("fc2_k5_funnel", int(small_m and self.trellis_bits == 5)),
+                ("grouped_fc2_grid376", int(not small_m)),
+                ("tile_major_tasks", int(self.tile_major_tasks and small_m)),
+                ("fc1_pipeline_stages", self.fc1_pipeline_stages if small_m else 2),
+                ("fc1_warps", self.fc1_warps if small_m else 4),
+                ("fc1_a_swizzle_rotate", int(self.fc1_a_swizzle_rotate and small_m)),
+                ("fc1_warp_quant", int(self.fc1_warp_quant and small_m)),
+                ("fc1_exact_staging", int(self.fc1_exact_staging and small_m)),
+                ("fc1_broadcast_a", int(self.fc1_broadcast_a and small_m)),
+                ("materialized", int(materialized)),
+                ("small_m_scheduler", int(small_m)),
+                ("fc1_tile_n", self.fc1_tile_n if small_m else 128),
+                ("experts", self.experts),
+                ("hidden", self.hidden),
+                ("intermediate", self.intermediate),
+                ("topk", self.topk),
+                ("rank", self.tp_rank),
+                ("scaled", 1),
+                ("identity", int(not self.full_coupled)),
+                ("scale_sandwich", int(self.scale_component is not None)),
+                ("full_coupled", int(self.full_coupled)),
+                ("raw_fc1_diagnostic", int(self.diagnostic_raw_fc1)),
+                ("input_prequant_diagnostic", int(self.diagnostic_raw_fc1)),
+                ("codebook", "mcg"),
+                ("deterministic_output", int(self.deterministic_output)),
+                ("down_remainder", int(self.p8_down_remainder)),
+                ("down_remainder_phases", self._dx2_phases if self.p8_down_remainder else "none"),
+                ("dx2_rowpack", int(self.p8_dx2_rowpack and small_m)),
+                ("dx2_input_rowpack", int(self.p8_input_rowpack and small_m)),
+                ("epi_par", int(self.p8_epi_par and small_m)),
+            ),
+            dsl_compile_options=OptLevel(2),
+        )
+        arm = _CompiledArm(compiled=compiled, tile_m=tile_m, materialized=materialized, mac=mac)
+        self._compiled[cache_key] = arm
+        return arm
+
+    def _compile_full_coupled_reducer(self):
+        if not self.full_coupled:
+            raise RuntimeError("coupled reducer requested for non-coupled P8")
+        if self._coupled_reducer is not None:
+            return self._coupled_reducer
+        from b12x.moe._shared.kernels.p8_coupled_topk import (
+            P8CoupledTopKSumKernel,
+        )
+
+        reducer = P8CoupledTopKSumKernel(topk=self.topk, hidden=self.hidden)
+        self._coupled_reducer = b12x_compile(
+            reducer,
+            make_ptr(cutlass.Float32, 16, cute.AddressSpace.gmem, assumed_align=16),
+            make_ptr(cutlass.Float32, 4, cute.AddressSpace.gmem, assumed_align=4),
+            make_ptr(cutlass.BFloat16, 16, cute.AddressSpace.gmem, assumed_align=16),
+            1,
+            current_cuda_stream(),
+            compile_spec=KernelCompileSpec.from_fields(
+                "glm53.p8.coupled_topk_h512",
+                2,
+                ("topk", self.topk),
+                ("hidden", self.hidden),
+                ("rank", self.tp_rank),
+                ("route_dtype", "fp32"),
+                ("output_dtype", "bf16"),
+            ),
+            dsl_compile_options=OptLevel(2),
+        )
+        return self._coupled_reducer
+
+    @torch.inference_mode()
+    def __call__(
+        self,
+        x: torch.Tensor,
+        topk_weights: torch.Tensor,
+        topk_ids: torch.Tensor,
+    ) -> torch.Tensor:
+        if x.dtype != torch.bfloat16 or x.ndim != 2 or x.shape[1] != self.hidden:
+            raise RuntimeError(f"P8 native input contract mismatch: {x.dtype} {tuple(x.shape)}")
+        m = int(x.shape[0])
+        if tuple(topk_ids.shape) != (m, self.topk) or tuple(topk_weights.shape) != (m, self.topk):
+            raise RuntimeError("P8 native routing shape mismatch")
+        if self.prefill_chunk_tokens and m > self.prefill_chunk_tokens:
+            if not self.full_coupled or self.debug_capture:
+                raise RuntimeError("bounded prefill requires full coupling without debug capture")
+            # MoE is token-local. Bound route-output and materialized carrier
+            # storage without changing the scheduler batch or attention work.
+            # Reuse allocations on the same stream; no host readback or sync.
+            output = torch.empty_like(x)
+            for start in range(0, m, self.prefill_chunk_tokens):
+                stop = min(start + self.prefill_chunk_tokens, m)
+                output[start:stop].copy_(self(x[start:stop], topk_weights[start:stop], topk_ids[start:stop]))
+            return output
+        # Every stored rate now has a fused grouped M64/N128 owner, so a non-K4 layer at M>1
+        # runs the same native path as K4 rather than looping the M1 kernel row by row. The
+        # row-by-row fallback is deliberately gone: a rate without a grouped specialization
+        # must fail closed instead of silently serving at a fraction of the speed.
+        if self.scale_component is not None and not self.full_coupled and m != 1:
+            raise RuntimeError("P8 scale sandwich currently supports M=1 only")
+        # Match the W4A8 planner's measured M16-to-M64 transition: sparse
+        # decode and ordinary prefill stay monolithic; only dense routed
+        # batches pay for the split materialized phase kernels.
+        materialized = (
+            m * self.topk >= 36 * self.experts
+            if self.force_materialized is None
+            else self.force_materialized
+        )
+        small_m = use_small_m(self.small_m_scheduler, m)
+        if self.full_coupled:
+            # Experimental direct-route batches through M16; requires matching
+            # multirow FC1 and input-prologue patches. Not serving-qualified.
+            small_m = m <= 16 and not self.grouped_m16
+            materialized = True
+        materialized = materialized or small_m
+        arm = self._compile(materialized, small_m=small_m, expected_m=m)
+        tile_m = arm.tile_m
+        x = x.contiguous()
+        flat_ids = topk_ids.to(dtype=torch.int32).contiguous().reshape(-1)
+        flat_weights = topk_weights.to(dtype=torch.float32).contiguous().reshape(-1)
+        physical_tiles = (
+            m * self.topk if small_m
+            else self.experts + (m * self.topk + tile_m - 1) // tile_m
+        )
+        rows_padded = physical_tiles * tile_m
+        gate_tile_count = ((2 * self.intermediate) // 128) // 2
+        max_tasks = physical_tiles * max(gate_tile_count, 1)
+        fused_scratch_zero = ((self.fuse_scratch_zero and small_m) or
+                              (self.fuse_grouped_scratch and self.grouped_m16))
+        shared_kernel_output = None
+        if fused_scratch_zero or self.shared_workspace:
+            if fused_scratch_zero:
+                from .direct_policy_scratch import direct_scratch_layout
+                self._scratch_layout = (direct_scratch_layout(m, self.intermediate, tile_m=tile_m, planes=self._dx2_planes) if small_m
+                    else p8_small_m_scratch_layout(intermediate=self.intermediate, tokens=m, planes=self._dx2_planes,
+                        shared=True, grouped=True, tile_m=tile_m, direct=False))
+            layout = (p8_small_m_scratch_layout(intermediate=self.intermediate, tokens=m, shared=True, planes=self._dx2_planes,
+                          grouped=self.full_coupled and materialized and not small_m, tile_m=tile_m, direct=small_m)
+                      if self.shared_workspace else self._scratch_layout)
+            assert layout is not None
+            # A single GPU fill initializes all original bytes plus alignment
+            # padding. The views add no casts, copies, or device kernels.
+            if self.shared_workspace:
+                from vllm.v1.worker.workspace import current_workspace_manager
+                arena, shared_kernel_output = current_workspace_manager().get_simultaneous(
+                    ((layout.nbytes,), torch.uint8),
+                    ((m * self.topk, self.hidden), torch.float32),
+                )
+                arena.zero_()
+            else:
+                arena = torch.zeros(layout.nbytes, dtype=torch.uint8, device=self.device)
+            buffers = {
+                region.name: arena.narrow(0, region.offset, region.nbytes)
+                .view(getattr(torch, region.dtype)).reshape(region.shape)
+                for region in layout.regions
+            }
+            packed_a = buffers["packed_a"]
+            scale_flat = buffers["scale_flat"]
+            intermediate_u32 = buffers["intermediate_u32"]
+            barrier_count = buffers["barrier_count"]
+            barrier_epoch = buffers["barrier_epoch"]
+            pair_head = buffers["pair_head"]
+            producers_done = buffers["producers_done"]
+            all_published = buffers["all_published"]
+            task_head = buffers["task_head"]
+            task_tail = buffers["task_tail"]
+            task_ready = buffers["task_ready"]
+            task_expert = buffers["task_expert"]
+            task_m_tile = buffers["task_m_tile"]
+            task_slice_begin = buffers["task_slice_begin"]
+            task_slice_count = buffers["task_slice_count"]
+            task_valid_rows = buffers["task_valid_rows"]
+            tile_write_count = buffers["tile_write_count"]
+            row_counts = buffers["row_counts"]
+            expert_write_rows = buffers["expert_write_rows"]
+            expert_tile_base = buffers["expert_tile_base"]
+            token_map = buffers["token_map"]
+            token_weights = buffers["token_weights"]
+            output = (torch.zeros(m, self.hidden, dtype=torch.bfloat16, device=self.device)
+                      if self.shared_workspace else buffers["output"])
+        else:
+            # Coupled grouped FC1 addresses the shared carrier by token index,
+            # not padded expert row (p8_h128_fc1 src_word uses tok). Keep the
+            # original extent for every other owner, including M1.
+            input_rows = (m if self.compact_input_storage and self.full_coupled
+                          and materialized and not small_m else rows_padded)
+            packed_a = torch.zeros(input_rows * self.hidden, dtype=torch.uint8, device=self.device)
+            scale_elements = (
+                m * (self.hidden // 32)
+                if self.full_coupled and materialized and not small_m
+                else (self.experts + m * self.topk + 1)
+                * tile_m
+                * (self.hidden // 8)
+            )
+            scale_flat = torch.zeros(
+                scale_elements, dtype=torch.uint8, device=self.device
+            )
+            intermediate_count = self._dx2_planes * rows_padded * (self.intermediate + self.intermediate // 32) // 4
+            intermediate_u32 = torch.zeros(intermediate_count, dtype=torch.int32, device=self.device)
+
+            def z1():
+                return torch.zeros(1, dtype=torch.int32, device=self.device)
+
+            def ztask():
+                return torch.zeros(max_tasks, dtype=torch.int32, device=self.device)
+
+            barrier_count, barrier_epoch = z1(), z1()
+            pair_head, producers_done, all_published = z1(), z1(), z1()
+            task_head, task_tail = z1(), z1()
+            task_ready, task_expert, task_m_tile = ztask(), ztask(), ztask()
+            task_slice_begin, task_slice_count, task_valid_rows = ztask(), ztask(), ztask()
+            tile_write_count = torch.zeros(physical_tiles, dtype=torch.int32, device=self.device)
+            row_counts = torch.zeros(self.experts, dtype=torch.int32, device=self.device)
+            expert_write_rows = torch.zeros(self.experts, dtype=torch.int32, device=self.device)
+            expert_tile_base = torch.zeros(self.experts + 1, dtype=torch.int32, device=self.device)
+            token_map = torch.zeros(rows_padded, dtype=torch.int32, device=self.device)
+            token_weights = torch.zeros(rows_padded, dtype=torch.float32, device=self.device)
+            output = torch.zeros(m, self.hidden, dtype=torch.bfloat16, device=self.device)
+        if self.diagnostic_raw_fc1:
+            # Keep the ordinary allocation block exactly unchanged. Only the
+            # diagnostic arm fills NaN payload sentinels and zeroes counters.
+            intermediate_u32.fill_(-1)
+            trace_base = rows_padded * (self.intermediate // 4)
+            intermediate_u32[trace_base + 32 : trace_base + 64].zero_()
+        kernel_output = (
+            shared_kernel_output if shared_kernel_output is not None else torch.empty(
+                m * self.topk,
+                self.hidden,
+                dtype=torch.float32 if self.full_coupled else torch.bfloat16,
+                device=self.device,
+            )
+            if self.deterministic_output
+            else output
+        )
+        launch_mac = arm.mac
+        if self.grid_policy and self.world_size == 4 and self.mac_override is None:
+            from .p8_multirow_scratch import direct_grid_capacity
+            launch_mac = direct_grid_capacity(m, arm.mac)
+        arm.compiled(
+            _gptr(cutlass.BFloat16, x),
+            _gptr(cutlass.Int32, flat_ids, 4),
+            _gptr(cutlass.Float32, flat_weights, 4),
+            _gptr(cutlass.Float4E2M1FN, packed_a),
+            _gptr(cutlass.Float8E4M3FN, scale_flat),
+            _gptr(cutlass.Uint8, packed_a),
+            _gptr(cutlass.Uint8, scale_flat),
+            _gptr(cutlass.Uint32, intermediate_u32),
+            barrier_count, barrier_epoch, pair_head, producers_done, all_published,
+            task_head, task_tail,
+            _gptr(cutlass.Int32, task_ready, 4),
+            _gptr(cutlass.Int32, task_expert, 4),
+            _gptr(cutlass.Int32, task_m_tile, 4),
+            _gptr(cutlass.Int32, task_slice_begin, 4),
+            _gptr(cutlass.Int32, task_slice_count, 4),
+            _gptr(cutlass.Int32, task_valid_rows, 4),
+            _gptr(cutlass.Int32, tile_write_count, 4),
+            self.w13_dummy,
+            _gptr(cutlass.Float8E4M3FN, self.sentinel),
+            self.w2_dummy,
+            _gptr(cutlass.Float8E4M3FN, self.sentinel),
+            _gptr(cutlass.Uint8, self.w13_scale_mx),
+            _gptr(cutlass.Uint8, self.w2_scale_mx),
+            _gptr(cutlass.Uint8, self.sentinel),
+            _gptr(cutlass.Uint8, self.sentinel),
+            _gptr(cutlass.Uint32, self.w13_stream),
+            _gptr(cutlass.Uint32, self.w13_sfb),
+            _gptr(cutlass.Uint32, self.w2_stream),
+            _gptr(cutlass.Uint32, self.w2_sfb),
+            row_counts, expert_write_rows, expert_tile_base,
+            self.ones, self.ones, self.ones, self.ones,
+            _gptr(
+                cutlass.Float32 if self.full_coupled else cutlass.BFloat16,
+                kernel_output,
+            ),
+            _gptr(cutlass.Int32, token_map, 4),
+            _gptr(cutlass.Float32, token_weights, 4),
+            m,
+            m * self.topk,
+            m * self.topk if self.deterministic_output else m,
+            rows_padded,
+            max_tasks,
+            physical_tiles,
+            launch_mac,
+            current_cuda_stream(),
+            _gptr(cutlass.Uint8, self.input_prequant_trace),
+            _gptr(cutlass.Float16, self.scale_component_packed),
+        )
+        if self.deterministic_output and not self.diagnostic_raw_fc1:
+            if self.full_coupled:
+                reducer = self._compile_full_coupled_reducer()
+                reducer(
+                    _gptr(cutlass.Float32, kernel_output),
+                    _gptr(cutlass.Float32, flat_weights, 4),
+                    _gptr(cutlass.BFloat16, output),
+                    m,
+                    current_cuda_stream(),
+                )
+            else:
+                _launch_dynamic_topk_sum(
+                    route_output=kernel_output,
+                    output=output,
+                    m=m,
+                    num_topk=self.topk,
+                    k=self.hidden,
+                    stream=current_cuda_stream(),
+                )
+        if self.debug_capture:
+            self.debug_tensors = {
+                "packed_a": packed_a, "scale_flat": scale_flat,
+                "intermediate_u32": intermediate_u32,
+                "route_output": kernel_output,
+                "token_map": token_map, "row_counts": row_counts,
+                "expert_tile_base": expert_tile_base,
+            }
+            self.debug_dispatch = {"small_m": small_m, "materialized": materialized,
+                                   "fused_scratch_zero": fused_scratch_zero,
+                                   "fc1_tile_n": self.fc1_tile_n if small_m else 128,
+                                   "tile_m": tile_m}
+            if self.diagnostic_raw_fc1:
+                self.debug_dispatch["diagnostic_raw_fc1"] = True
+                self.debug_tensors["input_prequant_trace"] = (
+                    self.input_prequant_trace
+                )
+        return output

runtime/patches/vllm_quant_trellismx.py

@@ -0,0 +1,195 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Opt-in native P8 routed experts over a ModelOpt NVFP4 carrier.
+
+The runtime is separately licensed and lazily imported. Dense, attention,
+router, shared-expert and MTP tensors retain the carrier's quantization.
+"""
+
+from vllm import envs
+
+import regex as re
+import torch
+
+from vllm.config import get_current_vllm_config
+from vllm.logger import init_logger
+from vllm.model_executor.layers.fused_moe.activation import MoEActivation
+from vllm.model_executor.layers.fused_moe.fused_moe_method_base import (
+    FusedMoEMethodBase,
+)
+from vllm.model_executor.layers.quantization.modelopt import ModelOptNvFp4FusedMoE
+from vllm.utils.b12x import B12xWarmupUnit
+from vllm.utils.trellismx import load_overlay, routed_layer
+
+logger = init_logger(__name__)
+
+
+def maybe_trellismx_method(config, layer, prefix):
+    directory = envs.VLLM_TRELLISMX_CHECKPOINT
+    index = routed_layer(prefix)
+    if not directory:
+        return None
+    text_config = get_current_vllm_config().model_config.hf_text_config
+    if text_config.model_type not in ("glm5_next_text", "glm5_next_mtp"):
+        raise ValueError(
+            "TrellisMX overlay currently requires the GLM5Next text adapter"
+        )
+    if index is None:
+        if re.fullmatch(
+            r"(?:model\.language_model|language_model\.model|model)"
+            r"\.layers\.45\.(?:mtp_block\.)?mlp\.experts",
+            prefix,
+        ):
+            return None
+        raise ValueError(f"Unrecognized TrellisMX routed-expert prefix: {prefix}")
+    if index == 45:
+        if (45, 0) not in load_overlay(directory).records:
+            return None
+        # MTP carrier metadata describes the replaced MXFP8 tensors. Use the
+        # established NVFP4 allocation ABI before loading native P8 sidecars.
+        config = getattr(config, "nvfp4_config", config)
+    if getattr(config, "quant_method", None) != "NVFP4":
+        raise ValueError("TrellisMX requires the pinned ModelOpt NVFP4 carrier")
+    return TrellisMXMoEMethod(config, layer.moe_config, directory, index)
+
+
+class TrellisMXMoEMethod(ModelOptNvFp4FusedMoE):
+    """Keep the carrier's weight-loader ABI; execute routed weights using P8."""
+
+    def __init__(self, config, moe_config, directory, layer_index):
+        # Do not select or compile an NVFP4 expert backend we never execute.
+        FusedMoEMethodBase.__init__(self, moe_config)
+        self.quant_config = config
+        self.use_a16 = False
+        self.use_global_sf = False
+        parallel = moe_config.moe_parallel_config
+        if (
+            parallel.tp_size != 4
+            or parallel.ep_size != 1
+            or moe_config.hidden_dim != 4096
+            or moe_config.intermediate_size_per_partition != 512
+            or moe_config.num_experts != 288
+            or moe_config.experts_per_token != 8
+            or moe_config.has_bias
+            or moe_config.is_lora_enabled
+            or moe_config.activation != MoEActivation.SILU
+            or moe_config.in_dtype != torch.bfloat16
+            or moe_config.swiglu_limit != 10.0
+        ):
+            raise ValueError(
+                "TrellisMX GLM adapter requires TP4, EP1 and GLM Flash shapes"
+            )
+        self.layer_index = layer_index
+        self.rank = parallel.tp_rank
+        self.overlay = load_overlay(directory)
+        self.runtime = None
+
+    @property
+    def is_monolithic(self):
+        return False
+
+    @property
+    def supports_eplb(self):
+        return False
+
+    def get_fused_moe_quant_config(self, layer):
+        return None
+
+    def process_weights_after_loading(self, layer):
+        if self.runtime is not None:
+            raise RuntimeError("TrellisMX hot weight replacement is unsupported")
+        from b12x.moe._shared.trellismx.p8_native_kernel import P8NativeTPMoE
+
+        device = layer.w13_weight.device
+        if device.type != "cuda" or torch.cuda.get_device_capability(device) != (12, 0):
+            raise ValueError("This TrellisMX runtime requires SM120 CUDA")
+        sidecar = self.overlay.sidecar(self.layer_index, self.rank)
+        record = self.overlay.records[self.layer_index, self.rank]
+        runtime = P8NativeTPMoE(
+            sidecar,
+            device=device,
+            tp_rank=self.rank,
+            world_size=4,
+            layer=self.layer_index,
+            expected_design_sha256=record["source_design_sha256"],
+            expected_transform_sha256=self.overlay.transform_hash,
+            topk=8,
+            hidden=4096,
+            intermediate=512,
+            swiglu_limit=10.0,
+            small_m_scheduler=True,
+            fc1_tile_n=128,
+            fuse_scratch_zero=True,
+            prefill_chunk_tokens=0,
+            grid_policy=True,
+            fc1_warp_quant=False,
+            fc1_broadcast_a=True,
+        )
+        # Release only replaced routed storage, after a successful load. Never
+        # touch the runner's router/shared experts or the separately owned MTP.
+        released = 0
+        for name in (
+            "w13_weight",
+            "w2_weight",
+            "w13_weight_scale",
+            "w2_weight_scale",
+            "w13_weight_scale_2",
+            "w2_weight_scale_2",
+            "w13_input_scale",
+            "w2_input_scale",
+        ):
+            value = getattr(layer, name)
+            released += value.numel() * value.element_size()
+            setattr(
+                layer,
+                name,
+                torch.nn.Parameter(
+                    torch.empty(0, dtype=value.dtype, device=value.device),
+                    requires_grad=False,
+                ),
+            )
+        self.runtime = runtime
+        layer.b12x_warmup_provider = self
+        logger.info(
+            "TrellisMX layer=%d rank=%d K%d E4M3/UE8M0-32 "
+            "coupled-h512-h128 released_carrier_bytes=%d",
+            self.layer_index,
+            self.rank,
+            record["bits"],
+            released,
+        )
+
+    def apply(
+        self, layer, x, topk_weights, topk_ids, shared_experts, shared_experts_input
+    ):
+        if self.runtime is None:
+            raise RuntimeError("TrellisMX routed weights were not loaded")
+        if x.shape[0] == 0:
+            return torch.empty_like(x)
+        return self.runtime(x, topk_weights, topk_ids)
+
+    def apply_monolithic(self, *args, **kwargs):
+        raise RuntimeError("TrellisMX routing belongs to the Jovian MoE runner")
+
+    def get_b12x_warmup_unit(self, layer, token_counts, output_dtype):
+        runtime = self.runtime
+        if runtime is None:
+            raise RuntimeError("TrellisMX warmup before weight load")
+
+        def compile():
+            for tokens in token_counts:
+                x = torch.zeros(
+                    (tokens, 4096), dtype=output_dtype, device=runtime.device
+                )
+                ids = torch.arange(8, dtype=torch.int32, device=runtime.device)
+                ids = ids.expand(tokens, 8).contiguous()
+                weights = torch.full((tokens, 8), 0.125, device=runtime.device)
+                runtime(x, weights, ids)
+
+        # Each layer owns compiled descriptors and scratch; do not deduplicate
+        # across layers solely because K/shape match.
+        return B12xWarmupUnit(
+            name="TrellisMX",
+            key=(type(self), self.layer_index, self.rank),
+            compile=compile,
+        )

runtime/patches/vllm_utils_trellismx.py

@@ -0,0 +1,130 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""CPU-only validation for the versioned TrellisMX routed overlay."""
+
+import hashlib
+import json
+import struct
+from dataclasses import dataclass
+from functools import lru_cache
+from pathlib import Path
+
+import regex as re
+
+_PREFIX = re.compile(
+    r"(?:model\.language_model|language_model\.model|model)"
+    r"\.layers\.(\d+)\.(?:mtp_block\.)?mlp\.experts"
+)
+
+
+def routed_layer(prefix: str) -> int | None:
+    match = _PREFIX.fullmatch(prefix)
+    if match is None:
+        return None
+    layer = int(match[1])
+    return layer if 3 <= layer <= 45 else None
+
+
+def sha256(path: Path) -> str:
+    digest = hashlib.sha256()
+    with path.open("rb") as source:
+        for chunk in iter(lambda: source.read(8 * 1024 * 1024), b""):
+            digest.update(chunk)
+    return digest.hexdigest()
+
+
+def _local_file(root: Path, relative: str) -> Path:
+    path = root / relative
+    # Symlinked local checkpoints are supported, but traversal in manifests is not.
+    if Path(relative).is_absolute() or ".." in Path(relative).parts:
+        raise ValueError(f"Unsafe TrellisMX path: {relative}")
+    if not path.is_file():
+        raise ValueError(f"Missing TrellisMX artifact: {path}")
+    return path
+
+
+@dataclass(frozen=True)
+class Overlay:
+    root: Path
+    records: dict
+    transform_hash: str
+
+    def sidecar(self, layer: int, rank: int, *, verify: bool = True) -> Path:
+        record = self.records[layer, rank]
+        path = _local_file(self.root, record["path"])
+        if verify and sha256(path) != record["sha256"]:
+            raise ValueError(f"TrellisMX weight hash mismatch: {path}")
+        with path.open("rb") as stream:
+            length_bytes = stream.read(8)
+            if len(length_bytes) != 8:
+                raise ValueError(f"Truncated safetensors header: {path}")
+            length = struct.unpack("<Q", length_bytes)[0]
+            if not 2 <= length <= 16 * 1024 * 1024:
+                raise ValueError(f"Invalid safetensors header size: {path}")
+            header = json.loads(stream.read(length))
+        metadata = header.get("__metadata__", {})
+        expected = {
+            "schema": "glm53-p8-coupled-h512-h128-tp4-rank.v1",
+            "layer": str(layer),
+            "rank": str(rank),
+            "world_size": "4",
+            "bits": str(record["bits"]),
+            "alphabet": "e4m3",
+            "scale": "ue8m0-k32",
+            "law": "procedural-mcg-alpha2",
+            "source_design_sha256": record["source_design_sha256"],
+        }
+        for name, value in expected.items():
+            if metadata.get(name) != value:
+                raise ValueError(f"TrellisMX {path}: incompatible {name}")
+        return path
+
+
+@lru_cache(maxsize=4)
+def load_overlay(directory: str) -> Overlay:
+    root = Path(directory).absolute()
+    manifest = json.loads(_local_file(root, "trellismx-manifest.json").read_text())
+    if manifest.get("schema") != "trellismx.hf-overlay-release.v1":
+        raise ValueError("Unsupported TrellisMX overlay schema")
+    if manifest.get("carrier") != {
+        "repo_id": "local-inference-lab/GLM-5.3-Flash-NVFP4",
+        "revision": "520de24eabf507659eaef7c70f14fd584527facc",
+    }:
+        raise ValueError("Unsupported TrellisMX carrier identity")
+    allocation = manifest.get("allocation", {})
+    expected_layers = {int(n) for n in allocation}
+    if expected_layers not in (set(range(3, 45)), set(range(3, 46))):
+        raise ValueError("GLM TrellisMX requires layers 3..44 and optional MTP45")
+    if any(type(bits) is not int or bits not in (4, 5) for bits in allocation.values()):
+        raise ValueError("This TrellisMX adapter supports K4/K5 only")
+    designs = {
+        sha256(_local_file(root, f"design/design-{index}.json")) for index in range(3)
+    }
+    transform = _local_file(root, "design/transform.json")
+    transform_config = json.loads(transform.read_text())
+    if (
+        transform_config.get("boundary") != "coupled-h512-h128-suh-svh-v1"
+        or transform_config.get("sign_draw") != 0
+        or transform_config.get("activation") != "silu-cap10"
+    ):
+        raise ValueError("Unsupported TrellisMX coupled transform")
+    records = {}
+    for record in manifest.get("files", []):
+        key = record["layer"], record["rank"]
+        if key in records or key[0] not in expected_layers or key[1] not in range(4):
+            raise ValueError("Duplicate or invalid TrellisMX layer/rank")
+        expected_path = f"sidecars/p8-layer-{key[0]:03d}-tp4-rank-{key[1]}.safetensors"
+        if record["path"] != expected_path:
+            raise ValueError("Unexpected TrellisMX sidecar path")
+        if record["bits"] != allocation[str(key[0])]:
+            raise ValueError("TrellisMX allocation and sidecar rate disagree")
+        if record["source_design_sha256"] not in designs:
+            raise ValueError("TrellisMX source design outside allowlist")
+        if not re.fullmatch(r"[0-9a-f]{64}", record["sha256"]):
+            raise ValueError("Invalid TrellisMX weight hash")
+        if _local_file(root, record["path"]).stat().st_size != record["bytes"]:
+            raise ValueError("TrellisMX sidecar size mismatch")
+        records[key] = record
+    if set(records) != {(layer, rank) for layer in expected_layers for rank in range(4)}:
+        raise ValueError("Incomplete TrellisMX TP4 inventory")
+    return Overlay(root, records, sha256(transform))

runtime/serve-codecv2-mtp.sh

@@ -0,0 +1,8 @@
+#!/bin/bash
+set -euo pipefail
+runtime_dir=$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")" && pwd)
+checkpoint_dir=$(dirname -- "$runtime_dir")
+python3 "$runtime_dir/install-mtp-runtime.py"
+export MODEL_ROOT=${MODEL_ROOT:-$checkpoint_dir/carrier}
+export VLLM_TRELLISMX_CHECKPOINT=${VLLM_TRELLISMX_CHECKPOINT:-$checkpoint_dir}
+exec /release/serve-rp2.sh "$@"

sidecars/p8-layer-045-tp4-rank-0.safetensors

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:1c699c04c08180eadad2f40a6c371295d1c46618db7029e7b7587e73a93c6dd6
+size 963497720

sidecars/p8-layer-045-tp4-rank-1.safetensors

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:4ac3d8d98b153f23d65c690cf0e24a98974ba52611642619544ad0311d3277c1
+size 963497720

sidecars/p8-layer-045-tp4-rank-2.safetensors

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:2750ae58989848196200dfa2991f40824ff6b52066f53319afd8dcc113c9bde6
+size 963497720

sidecars/p8-layer-045-tp4-rank-3.safetensors

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:c9071b28f75dfbbb6b16017721ddc185d2eb9ab021d81bdc1ad4d632cb846ddb
+size 963497720

trellismx-manifest.json

@@ -1,6 +1,6 @@
 {
   "schema": "trellismx.hf-overlay-release.v1",
-  "checkpoint": "codec-v2-uniform-k4",
+  "checkpoint": "codec-v2-uniform-k4-mtp45-k4",
   "packaging": "routed-overlay-requires-pinned-stock-carrier",
   "source_manifest_sha256": "dd4b391c015f9738cf7e4ed2c43b57cebc859f703e3502ca76ec3eb98d7397d8",
   "carrier": {
@@ -49,20 +49,18 @@
     "6": 4,
     "7": 4,
     "8": 4,
-    "9": 4
+    "9": 4,
+    "45": 4
   },
   "payload": {
-    "bytes_under_identity_files": -151878720,
-    "coupled_metadata_bytes": 151436544,
-    "file_bytes": 161867586048,
-    "identity_k4_file_bytes": 161715707328,
-    "logical_elements": 304405807104,
-    "safetensors_header_bytes": 564480,
-    "same_size_or_smaller": false,
-    "stored_bpw_including_metadata": 4.253994694462529,
+    "coupled_metadata_bytes": 155042176,
+    "file_bytes": 165721576928,
+    "logical_elements": 311653564416,
+    "safetensors_header_bytes": 578656,
+    "stored_bpw_including_metadata": 4.2539947133553015,
     "weight_payload_bpw": 4.25,
-    "weight_payload_bytes": 161715585024,
-    "note": "codec-v2 re-encode of every routed layer; see design/design-1.json"
+    "weight_payload_bytes": 165565956096,
+    "note": "codec-v2 re-encode of every routed layer; see design/design-1.json; MTP45 K4 included; per-module details in MTP-VERIFY-REPORT.json"
   },
   "files": [
     {
@@ -1576,7 +1574,62 @@
       "rank": 3,
       "bits": 4,
       "source_design_sha256": "82687b2cace99c681d01e84c5bef5cf5874091231072f03755b307f4a69b9c4d"
+    },
+    {
+      "layer": 45,
+      "rank": 0,
+      "bits": 4,
+      "path": "sidecars/p8-layer-045-tp4-rank-0.safetensors",
+      "bytes": 963497720,
+      "sha256": "1c699c04c08180eadad2f40a6c371295d1c46618db7029e7b7587e73a93c6dd6",
+      "source_design_sha256": "dc3f7e59c570d09f2dcf93921b4ce060bf36dfdcc1634c6cd615229bbc055d10"
+    },
+    {
+      "layer": 45,
+      "rank": 1,
+      "bits": 4,
+      "path": "sidecars/p8-layer-045-tp4-rank-1.safetensors",
+      "bytes": 963497720,
+      "sha256": "4ac3d8d98b153f23d65c690cf0e24a98974ba52611642619544ad0311d3277c1",
+      "source_design_sha256": "dc3f7e59c570d09f2dcf93921b4ce060bf36dfdcc1634c6cd615229bbc055d10"
+    },
+    {
+      "layer": 45,
+      "rank": 2,
+      "bits": 4,
+      "path": "sidecars/p8-layer-045-tp4-rank-2.safetensors",
+      "bytes": 963497720,
+      "sha256": "2750ae58989848196200dfa2991f40824ff6b52066f53319afd8dcc113c9bde6",
+      "source_design_sha256": "dc3f7e59c570d09f2dcf93921b4ce060bf36dfdcc1634c6cd615229bbc055d10"
+    },
+    {
+      "layer": 45,
+      "rank": 3,
+      "bits": 4,
+      "path": "sidecars/p8-layer-045-tp4-rank-3.safetensors",
+      "bytes": 963497720,
+      "sha256": "c9071b28f75dfbbb6b16017721ddc185d2eb9ab021d81bdc1ad4d632cb846ddb",
+      "source_design_sha256": "dc3f7e59c570d09f2dcf93921b4ce060bf36dfdcc1634c6cd615229bbc055d10"
     }
   ],
-  "runtime_image": "verdictai/trellismx@sha256:609a5fc1cd7d994ba32d9c03626c414d315947eb9f13fab474a15bc8dfbe0129"
+  "runtime_image": "verdictai/trellismx@sha256:609a5fc1cd7d994ba32d9c03626c414d315947eb9f13fab474a15bc8dfbe0129",
+  "mtp": {
+    "layer": 45,
+    "bits": 4,
+    "scale_policy": "signed-unit coupled vectors; CPU torch seed 530045; no prior r27 MTP scales",
+    "native_runtime_patch": "runtime/serve-codecv2-mtp.sh"
+  },
+  "main_model_payload": {
+    "bytes_under_identity_files": -151878720,
+    "coupled_metadata_bytes": 151436544,
+    "file_bytes": 161867586048,
+    "identity_k4_file_bytes": 161715707328,
+    "logical_elements": 304405807104,
+    "safetensors_header_bytes": 564480,
+    "same_size_or_smaller": false,
+    "stored_bpw_including_metadata": 4.253994694462529,
+    "weight_payload_bpw": 4.25,
+    "weight_payload_bytes": 161715585024,
+    "note": "codec-v2 re-encode of every routed layer; see design/design-1.json"
+  }
 }