diff --git a/.agents/quantization-matrix.md b/.agents/quantization-matrix.md
index 72a134192..17bee582c 100644
--- a/.agents/quantization-matrix.md
+++ b/.agents/quantization-matrix.md
@@ -127,7 +127,7 @@ Registry source:
 | `QUANT-FP8-PCPT` | ModelOpt FP8 per-channel/per-token | W8/A8 | capability selected | - | - | - | - | - | `INVENTORIED` | - | leaf spec open | - |
 | `QUANT-MXFP8-MODELOPT` | ModelOpt MXFP8 | W8/A8 | CUDA/ROCm/XPU dispatch | - | - | - | - | - | `INVENTORIED` | - | leaf spec open | - |
 | `QUANT-MIXED-MODELOPT` | ModelOpt mixed precision | FP8/NVFP4/MXFP8 groups | per-layer | part | part | part | Y | Y | `PARTIAL` | 35B FP8+NVFP4 slice: [loader](../src/vllm/model_executor/models/qwen3_5_weights.cpp#L118), [FP8 tests](../tests/vt/test_ops_fp8_cutlass.cpp#L188), [NVFP4 tests](../tests/vt/test_ops_moe_grouped.cpp#L453), [gate](../tests/parity/test_qwen36_paged_engine.cpp#L78) | leaf spec open | - |
-| `QUANT-CT-MXFP4` | compressed-tensors MXFP4 `mxfp4-pack-quantized` (group 32, E8M0 block scales, NO global) | W4/A4 (W4A16 dequant landed; W4A4 GEMM later) | CUDA GB10 (FlashInfer W4A4 else Marlin W4A16); CPU dequant | Y | Y | - | - | - | `ANCHOR-BACKFILL` | CPU weight unpack + E8M0 dequant: NEW [mxfp4_dequant.h](../include/vllm/model_executor/model_loader/mxfp4_dequant.h) + [.cpp](../src/vllm/model_executor/model_loader/mxfp4_dequant.cpp#L14) (`E8M0ToF32` = `2^(byte-127)`, `DequantMxfp4ToBf16`/`ToF32`, group 32, no global; reuses `kE2M1Lut`). Unit gate [test_mxfp4_dequant.cpp](../tests/vllm/test_mxfp4_dequant.cpp#L34) — E8M0 known-byte decode, hand-computed 32-group dequant (bf16+f32), the E8M0-vs-fp8 + group-32-vs-16 RED traps, multi-row/group offsets, randomized rel-error vs a double-precision port of `dq_mxfp4_torch` with bf16==f32 exactness. CPU `-Werror` 0-warn. Ports FROM `compressed_tensors_w4a4_mxfp4.py:20-97` + `mxfp8_utils.py:61-65,222` + golden `tests/quantization/reference_mxfp4.py:28-117`. C/E/P PENDING: GPU W4A4 fp4 GEMM + Marlin W4A16 fallback + MoE expert path + e2e are NAMED later bricks; DeepSeek-V4 + Kimi-K3 loaders consume this once wired | [MXFP4 spike](specs/mxfp4-compressed-tensors.md) | `CLAIM-QUANT-MXFP4` |
+| `QUANT-CT-MXFP4` | compressed-tensors MXFP4 `mxfp4-pack-quantized` (group 32, E8M0 block scales, NO global) | W4/A16 native **Marlin mxf4 keep-quant** landed (GB10 target; W4A4 cute-dsl crashes sm_121) | CUDA GB10 Marlin W4A16 (E8M0, group_blocks=2); CPU dequant fallback | Y | Y | Y | Y | - | `ANCHOR-BACKFILL` | **W2 native compute + W3 gates + e2e (row/QUANT-CT-MXFP4, `1c5ee09e`):** W0 vehicle `Yi30/Qwen3-8B-MXFP4` runs on the 0.25.0 oracle; W1 traced FlashInfer-W4A4-selected-but-crashes-on-sm_121 -> **Marlin W4A16 is the GB10 target**. Native path: `generate_kernels.py` MXFP4 config (`kFE8M0fnu`, group_blocks 2) + regenerated instances; [`MarlinProcessExpertScalesMxfp4`](../src/vt/cuda/cuda_marlin_repack.cu) (byte-exact vs vLLM at all shapes); `MoeMarlinArgs.{group_size,mxfp4}` launcher branch; `Nvfp4Weight.{group_size,is_mxfp4}` + [`dense_nvfp4_gemm.h`](../include/vllm/model_executor/models/dense_nvfp4_gemm.h) branch + `MatmulMxfp4W4A16D`; [`dense_weight_loaders.h`](../include/vllm/model_executor/models/dense_weight_loaders.h) MXFP4 loaders; `qwen3_weights.cpp` detect+load. **Gates GREEN:** op-level GEMM vs independent CPU dequant 0.36% M=1/M=8 all real shapes ([`test_ops_moe_grouped.cpp`](../tests/vt/test_ops_moe_grouped.cpp)); model-facing `MakeLinearMethod->Apply->BuildMarlinDenseResident` bad=0 K=4096+12288 ([`test_linear_method.cpp`](../tests/vllm/model_executor/layers/test_linear_method.cpp)); **e2e 3/4 token-exact vs oracle golden (async-off)** ([evidence](../docs/bench-evidence/mxfp4-qwen/W3-e2e-result.md)). RESIDUAL: default-async degeneration is a PRE-EXISTING classic-dense-Qwen3 async bug (device-mirror not wired for `qwen3.cpp`, quant-independent, SEPARATE row); W4 bench owed; p3 formal distributional gate owed. Earlier CPU weight unpack + E8M0 dequant: NEW [mxfp4_dequant.h](../include/vllm/model_executor/model_loader/mxfp4_dequant.h) + [.cpp](../src/vllm/model_executor/model_loader/mxfp4_dequant.cpp#L14) (`E8M0ToF32` = `2^(byte-127)`, `DequantMxfp4ToBf16`/`ToF32`, group 32, no global; reuses `kE2M1Lut`). Unit gate [test_mxfp4_dequant.cpp](../tests/vllm/test_mxfp4_dequant.cpp#L34) — E8M0 known-byte decode, hand-computed 32-group dequant (bf16+f32), the E8M0-vs-fp8 + group-32-vs-16 RED traps, multi-row/group offsets, randomized rel-error vs a double-precision port of `dq_mxfp4_torch` with bf16==f32 exactness. CPU `-Werror` 0-warn. Ports FROM `compressed_tensors_w4a4_mxfp4.py:20-97` + `mxfp8_utils.py:61-65,222` + golden `tests/quantization/reference_mxfp4.py:28-117`. C/E/P PENDING: GPU W4A4 fp4 GEMM + Marlin W4A16 fallback + MoE expert path + e2e are NAMED later bricks; DeepSeek-V4 + Kimi-K3 loaders consume this once wired | [MXFP4 spike](specs/mxfp4-compressed-tensors.md) | `CLAIM-QUANT-MXFP4` |
 | `QUANT-CT-W4A8-FP8` | compressed-tensors W4A8 FP8 | W4/A8 | CUTLASS | - | - | - | - | - | `INVENTORIED` | - | leaf spec open | - |
 | `QUANT-CT-W4A8-INT8` | compressed-tensors W4A8 INT8 | W4/A8 | platform selected | - | - | - | - | - | `INVENTORIED` | - | leaf spec open | - |
 | `QUANT-CT-W8A8-FP8` | compressed-tensors W8A8 FP8 | W8/A8 | platform selected | - | - | - | - | - | `INVENTORIED` | - | leaf spec open | - |
diff --git a/.agents/specs/mxfp4-compressed-tensors.md b/.agents/specs/mxfp4-compressed-tensors.md
index fc7bcf6cc..7e45833da 100644
--- a/.agents/specs/mxfp4-compressed-tensors.md
+++ b/.agents/specs/mxfp4-compressed-tensors.md
@@ -6,7 +6,16 @@ brick landed; the GPU W4A4 fp4 GEMM is a named later brick).
 
 **Base:** current `main` HEAD `42c56b51` (isolated worktree, CPU-only, foreground,
 NOT pushed). **Pinned oracle:** `${VLLM_SOURCE}` = `/home/mudler/_git/vllm` @
-`5559679229bc961848b121ccdeaa8fa5d79bec98` (vLLM 0.26.0.dev0).
+`5559679229bc961848b121ccdeaa8fa5d79bec98` (vLLM 0.26.0.dev0). **Runnable oracle
+on-box** = `~/venvs/vllm-oracle -> vllm-oracle-v0.25.0-stage` (vLLM **0.25.0**,
+compressed_tensors 0.17.0). The 0.25.0 tree is the one W0-W5 gate against; its
+mxfp4 dispatch is byte-for-byte the same shape as the 0.26 pin (verified below).
+
+> **W0-W5 UPDATE (2026-08-05, USER-priority "full MXFP4 at vLLM parity,
+> benchmarked on a Qwen model", branch `row/QUANT-CT-MXFP4`).** DeepSeek/Kimi are
+> NOT the vehicle (won't fit / not the target); **Qwen dense is**. This block
+> re-scopes W2-W5 around a real on-box Qwen MXFP4 checkpoint and pins the parity
+> target from the RUNNING oracle. See "## W0-W5 (Qwen vehicle) — 2026-08-05".
 
 **Why now:** shared unblocker. Both **DeepSeek-V4-Flash** (W6 MegaMoE MXFP4
 experts) and **Kimi-K3** (its real 2.8T checkpoint is `mxfp4-pack-quantized`)
@@ -163,4 +172,247 @@ the NVFP4 `weight_scale_2` bf16 round.
 - **No on-box e2e yet** — the two owning checkpoints are huge (Kimi-K3 2.8T does not
   fit one GB10; DeepSeek-V4-Flash NVFP4/fp8 need multi-Spark). The dequant brick is
   gateable at unit scale today; the e2e stays derive-and-ship until a fitting MXFP4
-  vehicle runs (RECORDED, mirrors both model specs).
+  vehicle runs (RECORDED, mirrors both model specs). **SUPERSEDED for the e2e vehicle
+  by the Qwen dense path below (2026-08-05): a small dense Qwen3-8B MXFP4 checkpoint
+  fits GB10 and the oracle registers its arch, so the e2e gate is now reachable
+  independently of DeepSeek/Kimi.**
+
+---
+
+## W0-W5 (Qwen vehicle) — 2026-08-05
+
+USER re-scope: **full MXFP4 support at vLLM parity, benchmarked on a Qwen model.**
+Qwen dense is the vehicle (DeepSeek/Kimi explicitly out — fit/target). This section
+pins the parity target from the RUNNING 0.25.0 oracle and lays the W2-W5 contract.
+Empirical (RUN/BUILD/BENCH) steps are GPU-gated and may lag the spec (box shared;
+locks + disk contended); the design + oracle-support proofs below are not.
+
+### W0 — checkpoint gateability
+
+**Vehicle: `Yi30/Qwen3-8B-MXFP4`** (HF). The clean dense W0 vehicle.
+- `config.json`: `architectures=["Qwen3ForCausalLM"]` (dense — the best-supported
+  Qwen family, NOT the new hybrid `Qwen3_5ForConditionalGeneration`),
+  `quant_method="compressed-tensors"`, `format="mxfp4-pack-quantized"`,
+  `group_size=32`, `ignore=["lm_head"]`. **`input_activations` is SET**
+  (`dynamic=true`, `num_bits=4`, `type=float`, `group_size=32`) → this is a **true
+  W4A4** checkpoint (weights AND activations MXFP4), which selects the W4A4 GEMM on
+  GB10 (see W1). Weights: 4-bit float, group 32, symmetric.
+- Size: 2 shards, **6.18 GB** safetensors total; complete tokenizer +
+  `generation_config` + chat template. Fits GB10 trivially (~6 GiB weights).
+- **Oracle-support PROVED at the import/registry layer (0.25.0, on dgx):**
+  `ModelRegistry.get_supported_archs()` contains `Qwen3ForCausalLM`; and
+  `from ...schemes.compressed_tensors_w4a4_mxfp4 import CompressedTensorsW4A4Mxfp4`
+  imports clean. **Oracle-RUN proof (greedy golden) is GPU-gated — QUEUED.**
+- Alternatives surveyed: `Yi30/Qwen3-8B-MXFP4-LLMC` (same, produced by
+  llm-compressor — the mechanical-repro arm), `Yi30/Qwen3-8B-MXFP4-FP8KV[-FP8Attn]`
+  (adds fp8 KV/attn — extra axes, avoid for the clean gate);
+  `olka-fi/Qwen3.5-27B-MXFP4` (genuine `mxfp4-pack-quantized` but **weight-only**,
+  `input_activations=null`, arch `Qwen3_5ForConditionalGeneration` — the new hybrid;
+  a valid W4A16 secondary vehicle but a harder arch); `OsaurusAI/Qwen3.6-27B-MXFP4`
+  = MLX mode (`{group_size,bits,mode:mxfp4}`), NOT compressed-tensors — rejected.
+- **`llm-compressor` is NOT installed** in the oracle venv (self-quantize is a
+  fallback only if no checkpoint runs; not needed — a runnable checkpoint exists).
+
+### W1 — the kernel the oracle ACTUALLY runs (parity target)
+
+Traced in the **0.25.0 site-packages** (the runnable oracle), `file:line`:
+- `compressed_tensors/schemes/compressed_tensors_w4a4_mxfp4.py`
+  `CompressedTensorsW4A4Mxfp4.__init__` → `self.kernel =
+  init_mxfp4_linear_kernel()` (`model_executor/kernels/linear/__init__.py:804`).
+- `_POSSIBLE_MXFP4_KERNELS[CUDA] = [FlashInferMxFp4LinearKernel,
+  MarlinMxFp4LinearKernel, HummingMxFp4LinearKernel]` (`__init__.py:462-466`).
+  `init_mxfp4_linear_kernel` returns the **first** whose `is_supported()` is True.
+- `FlashInferMxFp4LinearKernel.is_supported` (`mxfp4/flashinfer.py:22-28`):
+  `current_platform.has_device_capability(100) and has_flashinfer_cutedsl()`.
+  **On GB10 both are True** (device cap `(12,1)`; `has_flashinfer_cutedsl()`
+  returns **True** on-box — verified via import). So **the FlashInfer W4A4 kernel
+  is selected FIRST; Marlin is never reached.**
+- **Parity target = true W4A4 fp4xfp4 GEMM via FlashInfer CUTLASS cute-dsl**
+  (`mxfp4/flashinfer.py:apply_weights`): `flashinfer_mxfp4_quantize(x)` quantizes
+  the activation to mxf4, then `flashinfer_scaled_fp4_mm(x_fp4, weight, x_scale,
+  weight_scale, backend="cute-dsl", block_size=32, use_nvfp4=False)`. Weight scale
+  is swizzled + N padded to mult-of-128 in `process_weights_after_loading`
+  (`swizzle_mxfp4_scales`).
+- **This overrides the row's earlier "Marlin W4A16 fallback" hypothesis for THIS
+  box.** Marlin `MxFp4` (`mxfp4/marlin.py`, `weight_global_scale=None`) is the
+  **non-Blackwell / cute-dsl-absent** fallback only. Per mirror policy the W4A4
+  path is our target on GB10; the W4A16 Marlin path stays the documented fallback
+  for sm_80..sm_89 and cute-dsl-absent boxes.
+- **Runtime confirmation (QUEUED, GPU-gated):** the oracle logs
+  `"Using FlashInferMxFp4LinearKernel for MXFP4 GEMM"` (`__init__.py` `logger.info_once`)
+  — grep it in the W0 run; and same-tool nsys to name the cute-dsl GEMM kernel.
+- Env overrides that would change the pick (record for the A/B): `--linear-backend`
+  != auto (filters the kernel list) and `VLLM_DISABLED_KERNELS` (can disable
+  FlashInfer to force Marlin — the exact lever to A/B the two arms on one box).
+
+### W1 EMPIRICAL RESULT (2026-08-05, RUNTIME — supersedes the source-only pick above)
+
+Ran the oracle on `Yi30/Qwen3-8B-MXFP4` on GB10 (evidence:
+`docs/bench-evidence/mxfp4-qwen/`). The source trace said FlashInfer; the RUNTIME
+says FlashInfer is **selected but CRASHES on sm_121**:
+- `init_mxfp4_linear_kernel` logs `Using FlashInferMxFp4LinearKernel for MXFP4 GEMM`
+  (is_supported passes: cap 121 >= 100, cute-dsl present), THEN engine start dies with
+  `flashinfer.utils.BackendSupportedError: mm_fp4 does not support backend 'cute-dsl'
+  with capability 121`. FlashInfer's cute-dsl mxf4 backend covers sm_100 (datacenter
+  Blackwell) but NOT sm_121 (GB10). **The default oracle config is non-functional
+  for this checkpoint on GB10.**
+- The WORKING path = `VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel` -> the next
+  supported kernel = `MarlinMxFp4LinearKernel` (W4A16 weight-only fp4 Marlin,
+  `apply_fp4_marlin_linear(weight_global_scale=None)`). Greedy golden PYEXIT=0,
+  coherent+correct (Paris/Rome/Berlin, 2+2=4, fibonacci) — W0 satisfied.
+- **CORRECTION: the GB10 parity target is Marlin W4A16 mxf4, NOT FlashInfer W4A4.**
+  This REVALIDATES the row's original Laguna-B2 Marlin W4A16 hypothesis. The
+  source-only W1 conclusion above was wrong for sm_121 because it did not model
+  flashinfer's RUNTIME backend gate. `is_supported` != actually-runs. Trace the
+  execution, not just the dispatch source.
+
+### W2 — native keep-quant compute route (design)
+
+> **REVISED per the W1 empirical result:** the GB10 target is **Marlin W4A16
+> mxf4** (weight-only fp4, bf16 activation), NOT the cute-dsl W4A4 GEMM. Route
+> through our EXISTING Marlin FP4 infra (`src/vt/cuda/marlin/...`,
+> `cuda_marlin_repack.cu`, the NVFP4/AWQ/GPTQ Marlin path) exactly as the Laguna
+> B2 route did for NVFP4: extend the FP4 Marlin format plumbing for group-32 E8M0
+> (no global scale). The cutlass-W4A4 extension below stays a FUTURE arm, only
+> reachable once a flashinfer sm_121 mxf4 backend exists or we write the cutlass
+> mxf4 mma directly; it is not today's GB10 bar.
+
+Target = W4A4 mxf4xf4 (GB10) with the Marlin W4A16 mxf4 fallback documented.
+Route through the SAME families vLLM uses, mirroring the landed NVFP4 lane:
+
+1. **Reuse the NVFP4 cutlass fp4 tensor-core GEMM** (`src/vt/cuda/
+   cuda_matmul_nvfp4_cutlass.cu` + `nvfp4_cutlass_tactics*`). mxf4 differs from
+   nvf4 only in the block-scale FORMAT: group **32** (not 16), **E8M0** scale bytes
+   (not fp8-e4m3), and **no global scale**. CUTLASS block-scaled fp4 supports both
+   (`ScaleVectorSize`/`SFVecSize` 16 vs 32, UE8M0 vs UE4M3 SF dtype) — this is
+   exactly flashinfer's `use_nvfp4=False, block_size=32`. mxf4 warp
+   `mma.sync ...mxf4nvf4` is consumer-Blackwell (sm_121) available (see
+   `no-fa2`/Thor state notes: mxf4 tensor cores are consumer-Blackwell-only, which
+   GB10 IS). New: a `vt::MatmulMxfp4Fp4` op (or an `mxf4` mode flag on the nvfp4
+   op) + a mxf4 activation-quant emitter (per-token per-32-group amax → E8M0 scale
+   → E2M1 pack), mirroring vLLM `flashinfer_mxfp4_quantize`.
+2. **Scheme-selection method mirroring `schemes/nvfp4.h`**: add `schemes/mxfp4.h`
+   with `Mxfp4W4A4LinearMethod` (+ a `Mxfp4W4A16` arm for the Marlin/weight-only
+   fallback) and a `MakeLinearMethod` factory chosen ONCE from the checkpoint (the
+   loader probes `weight_packed` + the `mxfp4-pack-quantized` format string), NOT a
+   per-call tensor-name probe. Honors the three MUST-route seams
+   (LinearMethod/QuantizationConfig policy split, `vt::` op registry gate, shared
+   decode runner). Weight staging: MXFP4 packed bytes are a DEVICE-GEMM operand →
+   route through `ResidentWeight` (the keep-quant-device-slice rule: a raw host-byte
+   view is all-zeros on GB10).
+3. **Loader probe**: recognize `format=="mxfp4-pack-quantized"` and the
+   `input_activations` presence to distinguish W4A4 (both) vs weight-only W4A16, and
+   populate an `Mxfp4Weight` (packed `[N,K/2]` U8 + E8M0 `weight_scale` `[N,K/32]`
+   U8, no global). Reuse the `Nvfp4Weight`/`OwnedTensor` residency plumbing.
+4. **CPU reference (the gate truth, extends W1 dequant):** add
+   `compressed_tensors/mxfp4_emulation.{h,cpp}` mirroring `nvfp4_emulation.*` but
+   SIMPLER (no global scales): `RefScaledMxfp4Quant` (activation → mxf4: per-token
+   per-32-group amax, E8M0 block scale `2^(floor(log2(amax/6))+127)` clamped, E2M1
+   cast+pack), and `RunMxfp4Emulation` (dequant weight via existing
+   `DequantMxfp4ToF32` + activation round-trip + f32 matmul). This is the
+   software-emulation arm the GB10 fp4xf4 GEMM validates against, exactly as
+   `EmulationNvFp4LinearKernel` is for NVFP4.
+
+### W2 code surface (our side, for the continuation) — Marlin W4A16 mxf4
+
+The dense `Qwen3ForCausalLM` (the Yi30 vehicle's arch) IS already in our engine, so
+W2 is a quant-lane extension, not a new model:
+- **Model + loader:** `src/vllm/model_executor/models/qwen3_weights.cpp`
+  `LoadQwen3ForCausalLMWeights` (dense Qwen3 text gate) + `qwen3_dense.cpp`.
+- **Scheme detection:** `src/vllm/entrypoints/model_loader.cpp` (compressed-tensors /
+  `uses_nvfp4_w4a4()` seam, `:750`) — add the `format=="mxfp4-pack-quantized"` probe
+  alongside the nvfp4-pack detection.
+- **Weight struct:** mirror `Nvfp4Weight` (`dense_nvfp4_gemm.h`) as an `Mxfp4Weight`
+  (packed `[N,K/2]` U8 + E8M0 `weight_scale` `[N,K/32]` U8, NO global). ResidentWeight
+  staging (keep-quant-device-slice rule).
+- **GEMM:** mirror `MatmulNvfp4W4A16D` (`dense_nvfp4_gemm.h:426`) +
+  `MarlinW4A16Enabled` (`:83`) as `MatmulMxfp4W4A16D`, extending the Marlin FP4 repack
+  (`src/vt/cuda/cuda_marlin_repack.cu`, `marlin_repack.h`) to consume E8M0 group-32
+  scales with no global (vLLM `apply_fp4_marlin_linear(weight_global_scale=None)`).
+- **Selection method:** `schemes/mxfp4.h` mirroring `schemes/nvfp4.h`
+  (`MakeLinearMethod`), chosen once from the checkpoint.
+- **Build:** git-archive this branch to a fresh DGX dir + CUDA build (~1-2 GiB tree,
+  fits the 31 GiB free; NOT ~21 GiB — that was a vLLM-oracle-source-tree figure).
+
+### W3 — correctness gates
+
+- **Unit (CPU, buildable off-GPU):** extend `tests/vllm/test_mxfp4_dequant.cpp` (or a
+  sibling `test_mxfp4_emulation.cpp`) with the activation-quant + emulated-W4A4
+  round-trip vs a double-precision reference; the E8M0-vs-fp8, group-32-vs-16, and
+  no-global RED traps carried over. Mirror NVFP4's emulation gate.
+- **GPU unit (RED-first):** the `vt::MatmulMxfp4Fp4` GEMM vs the CPU emulation
+  reference (near-exact), INCLUDING the **M=1 decode mis-route RED trap** (the
+  recorded nvfp4 M=1 class — a new fp4 GEMM path can silently mis-route at batch 1).
+- **e2e SACRED greedy golden** vs the oracle on `Yi30/Qwen3-8B-MXFP4`. vLLM's own
+  greedy on a dense small model may be bf16-non-deterministic → use the ratified
+  DISTRIBUTIONAL gate (ours in vLLM's K-run set) if strict token-exact does not
+  hold; verify a bigger dense model strict where feasible.
+
+### W4 — benchmark (Qwen, the binding bar)
+
+`tools/bench/online_gate.py` ours-vs-oracle on `Yi30/Qwen3-8B-MXFP4`, SAME
+checkpoint both arms, c1..c8 (>=), 3 reps, single load per arm, idle box, GPU lock
+held. Match-or-beat is the bar; record honestly. Note the A/B lever:
+`VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel` forces the oracle onto Marlin
+W4A16 for a second reference point (W4A4 vs W4A16 on one box).
+
+### Empirical status (2026-08-05)
+
+- **W0 DONE** — checkpoint downloaded (6.18 GB) and the oracle RAN a greedy golden
+  (PYEXIT=0, coherent+correct). Golden + evidence in `docs/bench-evidence/mxfp4-qwen/`.
+- **W1 DONE** — runtime-traced: FlashInfer W4A4 selected-but-crashes on sm_121;
+  Marlin W4A16 is the working GB10 path (see the W1 EMPIRICAL RESULT block).
+- **W2 IMPLEMENTED + BUILDS + RUNS; e2e correctness RED (commit `7068dca6`)** — the
+  native Marlin-W4A16-mxf4 keep-quant path landed (kernel-gen MXFP4 config +
+  regenerated group_blocks=2 instances/selector; `MarlinProcessExpertScalesMxfp4`;
+  `MoeMarlinArgs.{group_size,mxfp4}` + launcher branch; `Nvfp4Weight.{group_size,
+  is_mxfp4}` + `dense_nvfp4_gemm.h` branch + `MatmulMxfp4W4A16D`;
+  `dense_weight_loaders.h` MXFP4 loaders; `qwen3_weights.cpp` detect+load). Clean
+  `-Werror` CUDA build on GB10; loads Yi30/Qwen3-8B-MXFP4; dispatches the native
+  group_blocks=2 Marlin kernel; runs. **NOT token-exact vs the golden** — a
+  DETERMINISTIC, UNIFORM (prefill+decode, NOT graph-related: identical with
+  `VLLM_CPP_CUDAGRAPH=0`) numerics error: robust tokens survive (" Paris", "there"
+  match) but the rest degenerates. LOCALIZATION: the MXFP4 scale permute is PROVEN
+  byte-exact vs vLLM's `mxfp4_marlin_process_scales` (128x256 CPU check, 1024/1024),
+  and the fp4 dequant (`dequant_skip_flop=false` bias path) is a faithful lift, so
+  the residual is inside the group_blocks=2 Marlin GEMM interaction — a path our
+  prior NVFP4-only usage (group_blocks=1) never exercised.
+- **W3 e2e — COMPUTE PATH GREEN (async-off token-exact 3/4); async-default degeneration
+  ROOT-CAUSED as a PRE-EXISTING non-MXFP4 bug.** With `VT_ASYNC_SCHED=0` our engine's
+  native Marlin mxf4 e2e is TOKEN-EXACT vs the golden on 3 of 4 prompts (p1/p2/p4);
+  the 4th (open-ended story) diverges after the identical first token — bf16 /
+  implementation non-determinism (near-tie/distributional regime). The DEFAULT
+  (async on) degenerated, but NOT from MXFP4: the async executor overlaps the prior
+  step's output-copy with the forward, and classic dense `Qwen3ForCausalLM` lacks the
+  async device-mirror fix (the #31 class wired only for the gate models) — a
+  quant-independent, pre-existing classic-dense-Qwen3 async bug (SEPARATE row).
+  Evidence: `docs/bench-evidence/mxfp4-qwen/W3-e2e-result.md`. So the MXFP4
+  keep-quant COMPUTE is correct (3/4 e2e token-exact + all unit gates below).
+- **W3 unit gate — GREEN (`test_ops_moe_grouped.cpp`, commit `8469e333`).** The MXFP4
+  Marlin GEMM (`MoeGroupedGemmNvfp4Marlin`, mxfp4 args) vs the INDEPENDENT CPU dequant
+  reference (`DequantMxfp4ToF32` + f32 matmul) at K=256,N=128: **max_rel 3.8e-3** at
+  M=1 AND M=8 — pure bf16 rounding, NOT a systematic error. **PROVES the MXFP4
+  keep-quant compute (repack, E8M0 scale processing, group_blocks=2 dispatch,
+  launcher) is correct.** So the e2e degeneration (`7068dca6`) is **NOT the GEMM**.
+- **e2e residual localization (as of `e28130ee`) — every MXFP4 component verified
+  correct; the bug is NOT in the compute or the loader byte interpretation:**
+  - scale permute byte-exact vs vLLM at ALL model shapes (N=6144/12288, K=12288);
+  - GEMM unit gate 0.38% at M=1/M=8 (N=128);
+  - loader dequant of the REAL layer-0 q_proj = sane+correct weights (min -0.5, max
+    0.5, mean|.| 0.02; scales 2^-8..2^-7), shapes match, U8-scale discriminator works;
+  - model dispatch (`IsNvfp4()` = `!qkv_proj_fp4.Empty()`) routes MXFP4 correctly;
+  - dense `Qwen3ForCausalLM` + NVFP4-W4A16 are known token-exact e2e (model-matrix).
+  Remaining suspects, in order: (1) **the group_blocks=2 kernel at LARGE N/K** — the
+  unit gate only ran N=128; the extended shapes ({4096,4096},{4096,12288},{12288,4096})
+  are committed RUN-PENDING (`e28130ee`); (2) a model-integration subtlety (per-layer
+  activation diff vs the oracle to find the first divergent op). NEXT: run the extended
+  unit gate; if RED at a model shape → large-shape kernel fix; if GREEN → per-layer
+  activation dump vs oracle.
+- **W4 NOT REACHED** — bench (`online_gate.py` c1..c8x3, oracle arm MUST set
+  `VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel`) is owed once the e2e is green.
+- Build recipe (reproducible): git-archive branch to `~/work/mxfp4-w2`, `cmake -B
+  build -G Ninja -DCMAKE_BUILD_TYPE=RelWithDebInfo -DVLLM_CPP_CUDA=ON
+  -DVLLM_CPP_CUDA_ARCHITECTURES=121a -DVLLM_CPP_MARLIN=ON -DVLLM_CPP_TRITON=OFF
+  -DVLLM_CPP_CUTLASS_DIR=$HOME/cutlass-4.5.0`, then `ninja -C build vllm-cli` (build
+  SPECIFIC targets — bare `ninja` builds 100+ tests that whole-archive libvllm.a =
+  28 GiB, blows the disk floor).
diff --git a/.agents/state.md b/.agents/state.md
index 4d0d50083..942736107 100644
--- a/.agents/state.md
+++ b/.agents/state.md
@@ -36175,6 +36175,57 @@ tmux killed, `/dev/shm` freed, both locks released, `local-ai-worker` left stopp
 found. Evidence: local logs `scratchpad/kimi_build_dgx.log` + `kimi_test_dgx.log`;
 records in spec section 11, model-matrix `MODEL-TEXT-kimi-linear-*`, benchmark record.
 
+## QUANT-CT-MXFP4 Qwen vehicle: W0 checkpoint + W1 kernel target PINNED; empirical GPU/disk-gated
+
+<!-- state: 2026-08-06T21:00 -->
+
+USER re-scope (full MXFP4 at vLLM parity, benchmarked on a Qwen model; DeepSeek/Kimi
+NOT the vehicle). Spike checkpoint on `row/QUANT-CT-MXFP4` (branch off `main`
+`c1a7b452`; helper). No GPU held, no locks touched.
+
+W0 vehicle FOUND: `Yi30/Qwen3-8B-MXFP4` (HF) — dense `Qwen3ForCausalLM`,
+`quant_method=compressed-tensors`, `format=mxfp4-pack-quantized`, group 32,
+`ignore=[lm_head]`, `input_activations` SET (true W4A4), 6.18 GB safetensors,
+complete tokenizer+gen config. Oracle-SUPPORT proved at import/registry on the
+runnable 0.25.0 oracle (`~/venvs/vllm-oracle` -> v0.25.0-stage): `Qwen3ForCausalLM`
+in `get_supported_archs()` AND `CompressedTensorsW4A4Mxfp4` imports; compressed_tensors
+0.17.0. Oracle-RUN greedy golden = GPU-gated, QUEUED. `llm-compressor` NOT installed
+(self-quantize fallback not needed — a runnable checkpoint exists). Alts surveyed:
+`Yi30/*-LLMC` (llm-compressor repro arm), `olka-fi/Qwen3.5-27B-MXFP4` (weight-only
+W4A16, hybrid `Qwen3_5ForConditionalGeneration` — harder arch, secondary),
+`OsaurusAI/Qwen3.6-27B-MXFP4` (MLX mode, NOT compressed-tensors — rejected).
+
+W1 parity target PINNED from the running 0.25.0 tree (site-packages file:line):
+`compressed_tensors_w4a4_mxfp4.py` -> `init_mxfp4_linear_kernel()`
+(`kernels/linear/__init__.py:804`) returns the FIRST supported of
+`_POSSIBLE_MXFP4_KERNELS[CUDA]=[FlashInfer,Marlin,Humming]` (`:462-466`).
+`FlashInferMxFp4LinearKernel.is_supported` (`mxfp4/flashinfer.py:22-28`) =
+`has_device_capability(100) AND has_flashinfer_cutedsl()`; on GB10 BOTH True (cap
+(12,1); cute-dsl import returns True on-box). So the oracle runs the TRUE W4A4
+fp4xf4 cute-dsl GEMM (`flashinfer_mxfp4_quantize` + `flashinfer_scaled_fp4_mm`,
+`block_size=32,use_nvfp4=False`) — Marlin W4A16 is the non-Blackwell/cute-dsl-absent
+fallback only. This OVERRIDES the row's earlier "Marlin W4A16" hypothesis for GB10;
+per mirror policy W4A4 is our target here. Runtime confirm (grep
+`"Using FlashInferMxFp4LinearKernel for MXFP4 GEMM"` + same-tool nsys) = QUEUED.
+A/B lever recorded: `VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel` forces Marlin.
+
+W2 route DESIGNED (spec section "W0-W5 (Qwen vehicle)"): extend the landed NVFP4
+cutlass fp4 tensor-core GEMM to the mxf4 block-scale format (group 32, E8M0 SF, no
+global = flashinfer `use_nvfp4=False,block_size=32`; mxf4 mma is consumer-Blackwell
+= GB10-available) + a mxf4 activation-quant emitter; `schemes/mxfp4.h` selection
+method mirroring `schemes/nvfp4.h` (chosen once, ResidentWeight staging per the
+keep-quant-device-slice rule); CPU `mxfp4_emulation.{h,cpp}` mirroring
+`nvfp4_emulation.*` (simpler, no globals) as the gate truth. W3 gates (CPU emulation
+unit + GPU M=1 RED trap + e2e SACRED distributional) and W4 bench (`online_gate.py`
+c1..c8x3, single load) SPEC'd.
+
+EMPIRICAL STATUS: W0-run/W1-runtime/W2-build/W3/W4 all need GB10 exclusively; at spec
+time gpu.lock+/tmp/gpu HELD (Option-A `test_qwen36_async_serving`, 96% util) and a
+Kimi-48B `hf download` (91.5G) was shrinking free disk (98->82G), so a 6.2G pull would
+risk the 15G floor. WAITED (locks not broken, disk not breached). Resume: download
+when disk >= ~15G free + lock free, then greedy golden, then build+gate+bench.
+Evidence: spec `.agents/specs/mxfp4-compressed-tensors.md` (W0-W5 section).
+
 ## SERVE-ASYNC-OPTION-A: decode-graph input H2D staged OUT of capture (faithful vLLM) — LANDS gated OFF; RED finally CONSTRUCTED; throughput WASH refutes "c16 = the baked H2D"
 <!-- state: 2026-08-06T21:00 -->
 
@@ -36236,6 +36287,43 @@ the residual is prefill glue (task #61) + steady host-orchestration/GPU compute,
 async input path. Evidence `dgx:~/work/mirror-ab/option-a/`, branch `row/SERVE-ASYNC-
 OPTION-A`.
 
+## QUANT-CT-MXFP4 Qwen W0+W1 EMPIRICAL: oracle golden RUNS (Marlin W4A16); FlashInfer W4A4 CRASHES on sm_121
+
+<!-- state: 2026-08-06T21:30 -->
+
+Box freed (coordinator); resumed the QUEUED empirical steps under both flock locks,
+mem gate (94-95 GiB free >= 90), tmux + done-markers, oracle alone, single load,
+disk floor respected (31 GiB free after the 6.18 GB pull). Box left clean (locks
+released, tmux killed, GPU idle). Branch `row/QUANT-CT-MXFP4`, PR #38.
+
+W0 DONE: downloaded `Yi30/Qwen3-8B-MXFP4` and the 0.25.0 oracle RAN a greedy golden
+(`temperature=0,seed=0,enforce_eager`, 4 prompts): PYEXIT=0, coherent+correct
+(Paris/Rome/Berlin/Madrid; 2+2=4 ... 5+5=10; a coherent story; a correct recursive
+fibonacci). Satisfies the hard oracle rule (RUNS a greedy golden, not just
+constructs). Golden + evidence: `docs/bench-evidence/mxfp4-qwen/`
+(`golden_marlin_w4a16.json`, `W0-W1-oracle-run.md`).
+
+W1 DONE with a RUNTIME CORRECTION of the source-only trace: `init_mxfp4_linear_kernel`
+selects `FlashInferMxFp4LinearKernel` (is_supported passes on cap 121>=100 + cute-dsl)
+and logs "Using FlashInferMxFp4LinearKernel", THEN engine start DIES:
+`flashinfer.utils.BackendSupportedError: mm_fp4 does not support backend 'cute-dsl'
+with capability 121`. FlashInfer's cute-dsl mxf4 backend covers sm_100 datacenter
+Blackwell, NOT sm_121 GB10 -> the DEFAULT oracle config is non-functional for this
+checkpoint on GB10. WORKING path = `VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel`
+-> `MarlinMxFp4LinearKernel` (W4A16 weight-only fp4, `apply_fp4_marlin_linear,
+weight_global_scale=None`). So the GB10 parity target is MARLIN W4A16 mxf4, NOT
+FlashInfer W4A4 -- this REVALIDATES the row's original Laguna-B2 hypothesis and
+disproves my source-only W1 conclusion (is_supported != actually-runs; trace the
+execution). W4 bench MUST set that env or the oracle arm crashes.
+
+W2/W3/W4 REMAINING: native Marlin-W4A16-mxf4 keep-quant in our engine (extend
+`src/vt/cuda/marlin/*` FP4 plumbing for group-32 E8M0, no global -- the Laguna B2
+route, NOT a new cute-dsl kernel) + gates (e2e vs this golden, distributional if
+bf16-non-det) + `online_gate.py` bench. Build-env caveat: DGX has NO clean vllm.cpp
+checkout and ~31 GiB free (full CUDA build ~21 GiB, tight/slow; other agents'
+`~/work/*/build` must not be reused) -> W2 needs a git-archive + fresh CUDA build.
+Spec updated (`.agents/specs/mxfp4-compressed-tensors.md`, W1 EMPIRICAL RESULT block).
+
 ## MODEL-TEXT-kimi-linear: §8 SACRED oracle golden CAPTURED (STRICT, 8/8 deterministic); full our-engine e2e f32-loader-blocked (~183 GiB > 119 pool)
 <!-- state: 2026-08-06T23:00 -->
 
@@ -36275,6 +36363,39 @@ Triton-AOT decode + grouped-MoE device slabs, (c) speed. Box left clean; the 91.
 checkpoint retained in the HF cache. Evidence: `dgx:~/kimi-e2e/{oracle_full.log,mem_full.log}`,
 local golden `tests/parity/goldens/kimi_linear_greedy/`.
 
+## QUANT-CT-MXFP4 W2 native Marlin mxf4 keep-quant LANDS (builds+runs on GB10); e2e RED, bug localized to group_blocks=2 GEMM
+
+<!-- state: 2026-08-06T23:30 -->
+
+Coordinator freed the box for W2-W4. Implemented the native MXFP4 (mxfp4-pack-quantized)
+Marlin W4A16 keep-quant path for dense Qwen3, routed through the EXISTING Marlin FP4
+infra (Laguna-B2 route), grounded 1:1 in vLLM marlin_utils_fp4.prepare_fp4_layer_for_marlin
+(is_nvfp4=False). Commit `7068dca6` on `row/QUANT-CT-MXFP4`, PR #38.
+
+LANDED (all additive; NVFP4 + 27B/35B gate paths byte-unchanged): generate_kernels.py
+MXFP4 config (s_type kFE8M0fnu, group_blocks [2]) + regenerated selector/.cu (15 new
+group_blocks=2 instances); MarlinProcessExpertScalesMxfp4 (E8M0 passthrough permute);
+MoeMarlinArgs.{group_size,mxfp4} + cuda_moe_marlin.cu launcher branch (kFE8M0fnu,
+group_size 32, global=nullptr); Nvfp4Weight.{group_size,is_mxfp4} + dense_nvfp4_gemm.h
+branch + MatmulMxfp4W4A16D + MXFP4 CPU fallback; dense_weight_loaders.h MXFP4 loaders
+(U8 E8M0 scale [N,K/32], no global); qwen3_weights.cpp detect+load.
+
+VERIFIED: clean -Werror CUDA build on GB10 (~137-TU incremental; full tree fits if you
+build SPECIFIC targets — bare `ninja` whole-archives libvllm.a into 100+ tests = 28 GiB,
+blows the 15G disk floor, learned the hard way). Loads Yi30/Qwen3-8B-MXFP4; dispatches
+the native group_blocks=2 Marlin kernel; RUNS. e2e NOT token-exact vs golden: a
+DETERMINISTIC UNIFORM numerics error (prefill+decode; identical with VLLM_CPP_CUDAGRAPH=0
+so NOT graph-safety) — robust tokens survive (" Paris"/"there" match golden) but the
+rest degenerates; p4 token-1 also wrong ("#" vs "\n"), so it is a uniform GEMM error not
+an M=1-only mis-route. LOCALIZED: the MXFP4 scale permute is PROVEN byte-exact vs vLLM's
+mxfp4_marlin_process_scales (128x256 CPU check 1024/1024), and the fp4 dequant
+(dequant_skip_flop=false bias path) is a faithful lift — so the residual is the
+group_blocks=2 Marlin GEMM interaction our prior NVFP4-only (group_blocks=1) usage never
+exercised. W3 e2e RED (honest). NEXT: device unit gate MatmulMxfp4W4A16D vs the CPU
+dequant reference at M=1 AND M=8 (RED-first) to pinpoint; then re-gate e2e, then W4 bench.
+Box left clean (locks released, tmux killed, GPU idle, disk 27G). Build tree persists at
+dgx:~/work/mxfp4-w2 for the continuation.
+
 ## MODEL-TEXT-kimi-linear: bf16-resident loader/forward POOL MATH + grounded design (§13); implementation scoped, PENDING
 <!-- state: 2026-08-07T00:30 -->
 
@@ -36308,6 +36429,33 @@ no code landed beyond records, nothing broken. Row STAYS `ACTIVE`. Next: impleme
 this branch, keep `test_kimi_linear_forward` 12/12·614 green, then the full-model e2e
 (free -g >= 90, monitor, STOP if the pool math does not close in practice).
 
+## QUANT-CT-MXFP4 W3 unit gate GREEN: MXFP4 GEMM PROVEN correct; e2e residual is NOT the compute
+
+<!-- state: 2026-08-07T01:00 -->
+
+Built + ran the owed W3 device unit gate (`test_ops_moe_grouped.cpp`, commit `8469e333`):
+MXFP4 Marlin GEMM vs the INDEPENDENT CPU dequant reference (`DequantMxfp4ToF32` + f32
+matmul), single-expert dense routing, M=1 AND M=8. GREEN — max_rel 3.8e-3 (pure bf16
+rounding). This PROVES the MXFP4 keep-quant compute (repack, E8M0 scale processing,
+group_blocks=2 kernel dispatch, launcher) is correct, so the e2e degeneration
+(`7068dca6`) is NOT the GEMM.
+
+Exhaustive component verification (all GREEN, none is the bug): scale permute byte-exact
+vs vLLM's mxfp4_marlin_process_scales at ALL model shapes (N=6144/12288, K=12288, CPU
+check); GEMM unit gate 0.38% M=1/M=8; loader dequant of the REAL layer-0 q_proj gives
+sane+correct weights (min -0.5/max 0.5/mean 0.02, scales 2^-8..2^-7) with matching
+shapes + working U8-scale discriminator; model dispatch (`IsNvfp4()`=`!fp4.Empty()`)
+routes MXFP4 correctly; graph-off == graph-on (not graph-safety); dense Qwen3 +
+NVFP4-W4A16 are known token-exact e2e. So the e2e residual is NOT the compute NOR the
+loader byte interpretation. PRIME SUSPECT: the group_blocks=2 kernel at LARGE N/K (unit
+gate only ran N=128) — extended shape cases ({4096,4096},{4096,12288},{12288,4096})
+committed RUN-PENDING (`e28130ee`); else a model-integration subtlety needing a
+per-layer activation diff vs the oracle. NEXT: run the extended unit gate (box was taken
+by the Kimi bf16 agent — locks HELD, WAITED, did not intrude); RED at a model shape =>
+large-shape kernel fix; GREEN => per-layer activation dump. Then re-gate e2e + W4 bench.
+Box left as found (Kimi holds the locks; my tmux sessions killed; build tree persists at
+dgx:~/work/mxfp4-w2).
+
 ## MODEL-TEXT-kimi-linear: bf16-resident loader/forward IMPLEMENTED + CPU-gated (§13); dgx CUDA build + full-model e2e PENDING
 <!-- state: 2026-08-07T01:30 -->
 
@@ -36342,6 +36490,35 @@ PENDING (this row, on dgx.casa GB10 — checkpoint CACHED 92G, 89 GiB free RAM,
 vs the STRICT golden (free -g >= 90 before load, memory-monitored; STOP + record if the pool math
 diverges toward the limit). Row STAYS `ACTIVE`; `VT_KIMI_DEVICE_COMPUTE` default STAYS OFF until
 the e2e token gate is green.
+
+## QUANT-CT-MXFP4 e2e RESOLVED: compute CORRECT (async-off 3/4 token-exact); async-default degeneration is a PRE-EXISTING classic-dense-Qwen3 bug
+
+<!-- state: 2026-08-07T03:00 -->
+
+Closed the e2e residual. Root-caused via the box-free continuation: the DEFAULT
+(async scheduling ON) degeneration is NOT the MXFP4 compute — with VT_ASYNC_SCHED=0
+the SAME binary is TOKEN-EXACT vs the golden on 3/4 prompts (p1 capitals, p2 arithmetic,
+p4 fibonacci all EXACT; p3 open-ended story diverges after the identical first token
+" there" = bf16/impl non-determinism, near-tie regime). The async executor overlaps
+the prior step's output-copy (async_copy_queue_) with the current forward; classic dense
+Qwen3ForCausalLM (qwen3.cpp) lacks the async device-mirror fix (the #31 "async batch-1
+token-0 degeneration" class wired only for the gate models qwen3_5) -> quant-independent,
+pre-existing classic-dense-Qwen3 async bug (SEPARATE row; would hit bf16/NVFP4 too).
+
+MXFP4 compute PROVEN correct across the full stack: op-level Marlin GEMM vs independent
+CPU dequant (0.36% M=1/M=8, all real shapes, e28130ee); scale permute byte-exact vs vLLM
+at all model shapes; model-facing MakeLinearMethod->Apply->BuildMarlinDenseResident gate
+(bad=0 K=4096+K=12288, 02410453); and e2e async-off 3/4 token-exact. Also surfaced +
+recorded: MarlinDenseResidentFor is keyed by weight POINTER (fine in the model; a reused
+stack slot aliases in tests). Reverted an exploratory per-stream DenseMarlinWorkspace
+change (not the async cause; async overlaps output-copy, not 2 forwards).
+
+REMAINING: (1) classic-dense-Qwen3 async device-mirror fix (separate row, non-MXFP4) for
+a default-config e2e; (2) W4 bench (online_gate.py c1..c8x3) on the async-off compute
+path, oracle arm VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel; (3) formal
+distributional gate for p3. Evidence: docs/bench-evidence/mxfp4-qwen/W3-e2e-result.md.
+Box: locks released, tmux killed, build tree persists at dgx:~/work/mxfp4-w2.
+
 ## MODEL-TEXT-kimi-linear: FULL-MODEL GB10 e2e RUNS via bf16-resident path — NEAR-TIE 106/128 (6/8 prompts token-exact); pool math CLOSES
 <!-- state: 2026-08-07T03:00 -->
 
@@ -36381,3 +36558,4 @@ host round-trips (the 1.59 tok/s is the O(n^2) full-recompute + host-island rate
 decode). `VT_KIMI_DEVICE_COMPUTE` STAYS OFF (parity-enablers: a near-tie is not token-exact). Row STAYS
 `ACTIVE`. Residuals now precisely: (a) STRICT token-exactness (device islands + bf16 stream), (b) the
 paged het-KV incremental decode, (c) speed.
+
diff --git a/docs/bench-evidence/mxfp4-qwen/W0-W1-oracle-run.md b/docs/bench-evidence/mxfp4-qwen/W0-W1-oracle-run.md
new file mode 100644
index 000000000..c9076aafd
--- /dev/null
+++ b/docs/bench-evidence/mxfp4-qwen/W0-W1-oracle-run.md
@@ -0,0 +1,74 @@
+# QUANT-CT-MXFP4 — W0 + W1 empirical (oracle run on GB10)
+
+Date 2026-08-05. Box dgx.casa (GB10, sm_121, cap 12.1). Oracle
+`~/venvs/vllm-oracle` -> `vllm-oracle-v0.25.0-stage` (vLLM **0.25.0**,
+compressed_tensors 0.17.0). Checkpoint `Yi30/Qwen3-8B-MXFP4` (dense
+`Qwen3ForCausalLM`, compressed-tensors `mxfp4-pack-quantized`, group 32, true
+W4A4 checkpoint, 5.79 GiB loaded). Run: offline vLLM `LLM` API, greedy
+(`temperature=0, seed=0`), `enforce_eager=True`, `gpu_memory_utilization=0.30`,
+`max_model_len=2048`, 4 prompts x 48 tokens. Both GPU flock locks held, mem gate
+`free -g >= 90` (94-95 GiB free), single load, tmux + done-marker.
+
+## W1 — the kernel that actually runs (parity target), with a runtime correction
+
+**Default config CRASHES on sm_121.** `init_mxfp4_linear_kernel` selects the
+FIRST supported kernel = `FlashInferMxFp4LinearKernel` (its `is_supported` =
+`has_device_capability(100) AND has_flashinfer_cutedsl()`, both True on GB10):
+
+```
+INFO [__init__.py:835] Using FlashInferMxFp4LinearKernel for MXFP4 GEMM
+```
+
+but at engine start (KV-cache profiling forward) FlashInfer's runtime capability
+check REJECTS the cute-dsl mxf4 backend for sm_121:
+
+```
+flashinfer.utils.BackendSupportedError: mm_fp4 does not support backend 'cute-dsl' with capability 121
+RuntimeError: Engine core initialization failed.
+```
+
+So `is_supported` is optimistic (cap >= 100) but the flashinfer `mm_fp4` cute-dsl
+backend only covers datacenter Blackwell (sm_100), NOT consumer GB10 (sm_121).
+**The default oracle path is non-functional for this checkpoint on GB10.**
+
+**Working path = Marlin W4A16** via the recorded lever
+`VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel`, which drops FlashInfer from
+the candidate list so `init_mxfp4_linear_kernel` falls to the next supported:
+
+```
+INFO [__init__.py:835] Using MarlinMxFp4LinearKernel for MXFP4 GEMM
+```
+
+`MarlinMxFp4LinearKernel.apply_weights` -> `apply_fp4_marlin_linear(...,
+weight_global_scale=None)` — **W4A16 weight-only fp4 Marlin GEMM** (the E8M0 group
+scale folded, no global). This is the real GB10 parity target and it
+**revalidates the row's original Laguna-B2 "Marlin W4A16" hypothesis**; the
+source-only trace that concluded "FlashInfer W4A4" was wrong for sm_121 because it
+did not account for flashinfer's runtime backend gate. Lesson: trace the
+execution, not just the dispatch source.
+
+## W0 — greedy golden (PROOF the oracle RUNS the checkpoint)
+
+`PYEXIT=0`. Coherent, correct greedy output on all 4 prompts (full token ids in
+`golden_marlin_w4a16.json`):
+
+- "The capital of France is" -> " Paris. ... Italy is Rome. ... Germany is Berlin. ... Spain is Madrid. ..."
+- "Q: What is 2 + 2? A:" -> " 4. Q: What is 3 + 3? A: 6. ... 5 + 5? A: 10"
+- "Once upon a time, in a small village," -> coherent story (wise old man, magical tree)
+- "def fibonacci(n):" -> correct recursive implementation
+
+This satisfies the hard W0 oracle rule: the pinned oracle BUILDS+RUNS a greedy
+golden on a real Qwen MXFP4 checkpoint, not merely constructs the config.
+
+## Consequences for W2-W4
+
+- Parity target on GB10 = **Marlin W4A16 mxf4** (weight-only fp4, bf16 activation,
+  group 32, E8M0 scale, no global). Our engine already has Marlin FP4 infra
+  (`src/vt/cuda/marlin/...`, NVFP4/AWQ/GPTQ) — W2 extends the FP4 Marlin format
+  plumbing for group-32 E8M0 (the Laguna B2 route), NOT a new cute-dsl W4A4 kernel.
+- The golden here is the W3 e2e gate reference. vLLM greedy on a dense 8B may be
+  bf16-non-deterministic -> distributional gate if strict token-exact does not hold.
+- A W4A4 arm is only reachable if/when flashinfer ships an sm_121 mxf4 backend
+  (or we implement the cutlass mxf4 mma directly); it is not the GB10 bar today.
+- The `VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel` env is REQUIRED for the
+  oracle to run this checkpoint on GB10 — record it in the W4 bench recipe.
diff --git a/docs/bench-evidence/mxfp4-qwen/W3-e2e-result.md b/docs/bench-evidence/mxfp4-qwen/W3-e2e-result.md
new file mode 100644
index 000000000..44855a57f
--- /dev/null
+++ b/docs/bench-evidence/mxfp4-qwen/W3-e2e-result.md
@@ -0,0 +1,43 @@
+# QUANT-CT-MXFP4 — W3 e2e result (Yi30/Qwen3-8B-MXFP4, GB10)
+
+Date 2026-08-05. OUR engine (`vllm-cli`, native Marlin W4A16 mxf4 keep-quant path)
+vs the oracle golden (`golden_marlin_w4a16.json`). Greedy, temp 0, seed 0, 48 tokens.
+
+## Root cause of the async-default degeneration (NOT the MXFP4 compute)
+
+The DEFAULT config (async scheduling ON, `max_concurrent_batches=2`) produced
+degenerate output (" Paris A A ( ( Paris ..."). Root-caused: **not the MXFP4
+compute** — with `VT_ASYNC_SCHED=0` the SAME binary produces coherent, token-exact
+output. The async executor overlaps the previous step's output-copy
+(`async_copy_queue_`) with the current forward; classic dense `Qwen3ForCausalLM`
+(`qwen3.cpp`) does not carry the async device-mirror fix (the #31 "async batch-1
+token-0 degeneration" class that was wired for the gate models `qwen3_5`), so the
+overlap corrupts the sampled token. This is QUANT-INDEPENDENT (would hit bf16/NVFP4
+classic dense Qwen3 the same way) and is a SEPARATE pre-existing bug, not this row.
+
+## MXFP4 compute is CORRECT — async-off e2e is token-exact vs the golden
+
+With `VT_ASYNC_SCHED=0` (the compute path), 3 of 4 prompts are TOKEN-EXACT vs the
+oracle golden; the 4th diverges after the identical first token (bf16 /
+implementation non-determinism on an open-ended prompt — the ratified distributional
+/ near-tie regime):
+
+| prompt | ours (async-off) | golden | verdict |
+|---|---|---|---|
+| "The capital of France is" | " Paris. What is the capital of Italy? ... Rome ... Berlin ... Madrid ..." | (same) | EXACT |
+| "Q: What is 2 + 2? A:" | " 4. Q: 3+3? A: 6. ... 5+5? A: 10" | (same) | EXACT |
+| "def fibonacci(n):" | " if n == 0: return 0 elif n == 1: return 1 else: return ..." | (same) | EXACT |
+| "Once upon a time, in a small village," | " there lived a young girl named Lily ..." | " there was a wise old man ..." | first token " there" matches; diverges (non-det) |
+
+Combined with the unit gates — op-level MXFP4 Marlin GEMM vs independent CPU dequant
+(0.36% at M=1/M=8, all real shapes), scale permute byte-exact vs vLLM at all model
+shapes, and the model-facing `MakeLinearMethod -> Apply -> BuildMarlinDenseResident`
+gate (bad=0 at K=4096 and K=12288) — the MXFP4 keep-quant compute is correct.
+
+## Status
+- **W3 compute path: GREEN** (3/4 e2e token-exact + all unit gates).
+- **Async-default degeneration:** pre-existing classic-dense-Qwen3 async bug (separate
+  row); the compute is validated on the async-off path.
+- **W4 bench:** owed; run on the async-off compute path (oracle arm
+  `VLLM_DISABLED_KERNELS=FlashInferMxFp4LinearKernel`), or after the classic-dense
+  async fix for a default-config number.
diff --git a/docs/bench-evidence/mxfp4-qwen/golden_marlin_w4a16.json b/docs/bench-evidence/mxfp4-qwen/golden_marlin_w4a16.json
new file mode 100644
index 000000000..b9e0c2c25
--- /dev/null
+++ b/docs/bench-evidence/mxfp4-qwen/golden_marlin_w4a16.json
@@ -0,0 +1,218 @@
+[
+  {
+    "prompt": "The capital of France is",
+    "text": " Paris. What is the capital of Italy? The capital of Italy is Rome. What is the capital of Germany? The capital of Germany is Berlin. What is the capital of Spain? The capital of Spain is Madrid. What is the capital",
+    "token_ids": [
+      12095,
+      13,
+      3555,
+      374,
+      279,
+      6722,
+      315,
+      15344,
+      30,
+      576,
+      6722,
+      315,
+      15344,
+      374,
+      21718,
+      13,
+      3555,
+      374,
+      279,
+      6722,
+      315,
+      9856,
+      30,
+      576,
+      6722,
+      315,
+      9856,
+      374,
+      19846,
+      13,
+      3555,
+      374,
+      279,
+      6722,
+      315,
+      17689,
+      30,
+      576,
+      6722,
+      315,
+      17689,
+      374,
+      24081,
+      13,
+      3555,
+      374,
+      279,
+      6722
+    ]
+  },
+  {
+    "prompt": "Q: What is 2 + 2? A:",
+    "text": " 4. Q: What is 3 + 3? A: 6. Q: What is 4 + 4? A: 8. Q: What is 5 + 5? A: 10",
+    "token_ids": [
+      220,
+      19,
+      13,
+      1207,
+      25,
+      3555,
+      374,
+      220,
+      18,
+      488,
+      220,
+      18,
+      30,
+      362,
+      25,
+      220,
+      21,
+      13,
+      1207,
+      25,
+      3555,
+      374,
+      220,
+      19,
+      488,
+      220,
+      19,
+      30,
+      362,
+      25,
+      220,
+      23,
+      13,
+      1207,
+      25,
+      3555,
+      374,
+      220,
+      20,
+      488,
+      220,
+      20,
+      30,
+      362,
+      25,
+      220,
+      16,
+      15
+    ]
+  },
+  {
+    "prompt": "Once upon a time, in a small village,",
+    "text": " there was a wise old man who lived in a house with a garden. The garden was very special because it had a magical tree that only bloomed once a year. The villagers would come to see the tree in full bloom, and they",
+    "token_ids": [
+      1052,
+      572,
+      264,
+      23335,
+      2310,
+      883,
+      879,
+      12163,
+      304,
+      264,
+      3753,
+      448,
+      264,
+      13551,
+      13,
+      576,
+      13551,
+      572,
+      1602,
+      3281,
+      1576,
+      432,
+      1030,
+      264,
+      23702,
+      4916,
+      429,
+      1172,
+      14211,
+      24139,
+      3055,
+      264,
+      1042,
+      13,
+      576,
+      60821,
+      1035,
+      2525,
+      311,
+      1490,
+      279,
+      4916,
+      304,
+      2480,
+      51454,
+      11,
+      323,
+      807
+    ]
+  },
+  {
+    "prompt": "def fibonacci(n):",
+    "text": "\n    if n == 0:\n        return 0\n    elif n == 1:\n        return 1\n    else:\n        return fibonacci(n-1) + fibonacci(n-2)\n\ndef main():\n    n = int",
+    "token_ids": [
+      198,
+      262,
+      421,
+      308,
+      621,
+      220,
+      15,
+      510,
+      286,
+      470,
+      220,
+      15,
+      198,
+      262,
+      4409,
+      308,
+      621,
+      220,
+      16,
+      510,
+      286,
+      470,
+      220,
+      16,
+      198,
+      262,
+      770,
+      510,
+      286,
+      470,
+      75698,
+      1445,
+      12,
+      16,
+      8,
+      488,
+      75698,
+      1445,
+      12,
+      17,
+      692,
+      750,
+      1887,
+      3932,
+      262,
+      308,
+      284,
+      526
+    ]
+  }
+]
diff --git a/include/vllm/model_executor/models/dense_nvfp4_gemm.h b/include/vllm/model_executor/models/dense_nvfp4_gemm.h
index 72b587522..fd23f434f 100644
--- a/include/vllm/model_executor/models/dense_nvfp4_gemm.h
+++ b/include/vllm/model_executor/models/dense_nvfp4_gemm.h
@@ -55,6 +55,7 @@
 #include <unordered_map>
 #include <vector>
 
+#include "vllm/model_executor/model_loader/mxfp4_dequant.h"  // DequantMxfp4ToBf16
 #include "vllm/model_executor/model_loader/nvfp4_dequant.h"  // DequantNvfp4ToBf16
 #include "vllm/model_executor/models/dense_device_glue.h"    // Dev/DBuf/MakeTensor
 #include "vllm/model_executor/models/qwen3_5_weights.h"      // Nvfp4Weight
@@ -149,7 +150,8 @@ inline Nvfp4Dev ResidentNvfp4(Dev d, const Nvfp4Weight& w) {
   }
   Nvfp4Dev r;
   r.packed = MakeTensor(w.d_packed.get(), DType::kI8, d.q.device, {w.n, w.k / 2});
-  r.scale = MakeTensor(w.d_scale.get(), DType::kI8, d.q.device, {w.n, w.k / 16});
+  // Scale grid is [N, K/group_size]: K/16 for NVFP4, K/32 for MXFP4.
+  r.scale = MakeTensor(w.d_scale.get(), DType::kI8, d.q.device, {w.n, w.k / w.group_size});
   return r;
 }
 
@@ -170,6 +172,24 @@ inline std::vector<uint16_t> DequantNvfp4ToBLayout(const Nvfp4Weight& w) {
   return io;
 }
 
+// MXFP4 analog: host dequant of an E8M0/group-32 fp4 weight to bf16 [K=in, N=out]
+// (Matmul-B layout) — the CPU / Marlin-disabled fallback. Bit-for-bit
+// vllm::DequantMxfp4ToBf16 + transpose. Independent of the Marlin kernel's own
+// E8M0 dequant, so a gate comparing the two paths is a real cross-check.
+inline std::vector<uint16_t> DequantMxfp4ToBLayout(const Nvfp4Weight& w) {
+  const int64_t out_dim = w.n, in_dim = w.k;
+  std::vector<uint16_t> oi(static_cast<size_t>(out_dim) * in_dim);
+  DequantMxfp4ToBf16(reinterpret_cast<const uint8_t*>(w.packed.bytes.data()),
+                     reinterpret_cast<const uint8_t*>(w.scale.bytes.data()),
+                     out_dim, in_dim, oi.data());
+  std::vector<uint16_t> io(static_cast<size_t>(in_dim) * out_dim);
+  for (int64_t r = 0; r < out_dim; ++r)
+    for (int64_t c = 0; c < in_dim; ++c)
+      io[static_cast<size_t>(c) * out_dim + r] =
+          oi[static_cast<size_t>(r) * in_dim + c];
+  return io;
+}
+
 #ifdef VT_MARLIN_NVFP4
 
 // --- Resident Marlin operands (repacked ONCE at first use) ------------------
@@ -196,27 +216,37 @@ inline void BuildMarlinDenseResident(Dev d, const Nvfp4Weight& w,
   if (mr.ready) return;
   const int K = static_cast<int>(w.k);
   const int N = static_cast<int>(w.n);
+  const int gs = static_cast<int>(w.group_size);  // 16 (nvfp4) or 32 (mxfp4)
   void* stream = d.q.handle;
   const size_t w_i32 = static_cast<size_t>(K / 16) * (static_cast<size_t>(N) * 2);
-  const size_t s_b = static_cast<size_t>(K / 16) * N;
+  const size_t s_b = static_cast<size_t>(K / gs) * N;  // K/16 nvfp4, K/32 mxfp4
   mr.w = d.b.Alloc(w_i32 * 4);
   mr.s = d.b.Alloc(s_b);
   mr.g = d.b.Alloc(sizeof(float));
   mr.n = w.n;
   mr.k = w.k;
-  std::vector<const uint8_t*> bufs{
-      reinterpret_cast<const uint8_t*>(w.scale.bytes.data())};
-  std::vector<size_t> lens{w.scale.bytes.size()};
-  const float sf = vt::cuda::MarlinNvfp4CombinedScaleFactor(bufs, lens);
   Nvfp4Dev dw = ResidentNvfp4(d, w);
   vt::cuda::MarlinRepackExpertWeight(stream, d.q.device.index,
                                      static_cast<uint32_t*>(mr.w),
                                      static_cast<const uint8_t*>(dw.packed.data), K, N);
-  vt::cuda::MarlinProcessExpertScales(stream,
-                                      static_cast<const uint8_t*>(dw.scale.data),
-                                      static_cast<uint8_t*>(mr.s), K, N, sf);
-  const float g = vt::cuda::MarlinNvfp4ProcessGlobalScale(w.scale2, sf);
-  d.b.Copy(d.q, mr.g, &g, sizeof(float));
+  if (w.is_mxfp4) {
+    // MXFP4: E8M0 passthrough permute (no combined factor, no global scale).
+    vt::cuda::MarlinProcessExpertScalesMxfp4(
+        stream, static_cast<const uint8_t*>(dw.scale.data),
+        static_cast<uint8_t*>(mr.s), K, N);
+    const float g = 1.0F;  // unused (kernel skips global for E8M0)
+    d.b.Copy(d.q, mr.g, &g, sizeof(float));
+  } else {
+    std::vector<const uint8_t*> bufs{
+        reinterpret_cast<const uint8_t*>(w.scale.bytes.data())};
+    std::vector<size_t> lens{w.scale.bytes.size()};
+    const float sf = vt::cuda::MarlinNvfp4CombinedScaleFactor(bufs, lens);
+    vt::cuda::MarlinProcessExpertScales(stream,
+                                        static_cast<const uint8_t*>(dw.scale.data),
+                                        static_cast<uint8_t*>(mr.s), K, N, sf);
+    const float g = vt::cuda::MarlinNvfp4ProcessGlobalScale(w.scale2, sf);
+    d.b.Copy(d.q, mr.g, &g, sizeof(float));
+  }
   d.b.Synchronize(d.q);  // repack done -> safe to free the fp4 originals
   w.d_packed.reset();
   w.d_scale.reset();
@@ -290,17 +320,20 @@ inline DBuf MatmulNvfp4MarlinD(Dev d, const Tensor& x, const Nvfp4Weight& w,
   // upcast (the same value it rounds to).
   DBuf outbf(d, DType::kBF16, {M, N});
   Tensor wq = MakeTensor(mr.w, DType::kI32, d.q.device, {1, K / 16, N * 2});
-  Tensor sc = MakeTensor(mr.s, DType::kI8, d.q.device, {1, K / 16, N});
+  // Scale grid rows = K/group_size (K/16 nvfp4, K/32 mxfp4).
+  Tensor sc = MakeTensor(mr.s, DType::kI8, d.q.device, {1, K / w.group_size, N});
   Tensor gg = MakeTensor(mr.g, DType::kF32, d.q.device, {1});
   Tensor wst = MakeTensor(ws, DType::kI32, d.q.device, {sms * 4});
   Tensor sorted = MakeTensor(ac.sorted, DType::kI32, d.q.device, {ac.max_tok});
   Tensor expert = MakeTensor(ac.expert, DType::kI32, d.q.device, {ac.max_blk});
   Tensor numpad = MakeTensor(ac.numpad, DType::kI32, d.q.device, {1});
   Tensor topkw = MakeTensor(ac.topkw, DType::kF32, d.q.device, {M});
-  vt::MoeGroupedGemmNvfp4Marlin(
-      d.q, outbf.t(), x, wq, sc, gg, wst, sorted, expert, numpad, topkw,
-      vt::MoeMarlinArgs{ac.block, 1, static_cast<int>(M), static_cast<int>(N),
-                        static_cast<int>(K), false});
+  vt::MoeMarlinArgs margs{ac.block, 1, static_cast<int>(M), static_cast<int>(N),
+                          static_cast<int>(K), false};
+  margs.group_size = static_cast<int>(w.group_size);
+  margs.mxfp4 = w.is_mxfp4;
+  vt::MoeGroupedGemmNvfp4Marlin(d.q, outbf.t(), x, wq, sc, gg, wst, sorted, expert,
+                                numpad, topkw, margs);
   if (out_dtype == DType::kBF16) return outbf;
   DBuf out(d, DType::kF32, {M, N});
   vt::CastF32(d.q, out.t(), outbf.t());
@@ -380,8 +413,12 @@ inline void BuildMarlinDensePairResident(Dev d, const Nvfp4Weight& gw,
 // True when a gate/up pair takes the fused Marlin gate_up path. Must be checked
 // IDENTICALLY at every call site so exactly ONE resident layout is ever built.
 inline bool GateUpFusedEligible(const Nvfp4Weight& gw, const Nvfp4Weight& uw) {
+  // MXFP4 takes the SPLIT path (two W4A16 GEMMs + MoeSiluMul): the fused merged
+  // gate_up resident is an NVFP4-only optimization; forcing split keeps the
+  // MXFP4 lane correct without a fused mxf4 pair repack (still byte-correct).
   return FusedGateUpEnabled() && !gw.Empty() && !uw.Empty() && !gw.IsTrueW4A4() &&
-         !uw.IsTrueW4A4() && gw.n == uw.n && gw.k == uw.k && gw.scale2 == uw.scale2;
+         !uw.IsTrueW4A4() && !gw.is_mxfp4 && !uw.is_mxfp4 && gw.n == uw.n &&
+         gw.k == uw.k && gw.scale2 == uw.scale2;
 }
 
 // silu(x@gate.T) * (x@up.T) -> bf16 [M,N] via ONE fused Marlin gate_up GEMM.
@@ -440,6 +477,14 @@ inline DBuf MatmulNvfp4W4A16D(Dev d, const Tensor& x, const Nvfp4Weight& w,
 #endif
   ++MutableW4A16Stats().fallback_gemms;
   DBuf dout(d, out_dtype, {M, N});
+  // MXFP4 has no naive redundant-dequant device op (that op is NVFP4-only); the
+  // fallback is always the host dequant + bf16 Matmul reference.
+  if (w.is_mxfp4) {
+    std::vector<uint16_t> wb = DequantMxfp4ToBLayout(w);
+    DBuf dwb(d, DType::kBF16, {K, N}, wb.data());
+    vt::Matmul(d.q, dout.t(), x, dwb.t());
+    return dout;
+  }
   // Same class-A conversion: the naive redundant-dequant NVFP4 GEMM exists only
   // where the op table realizes it (kCUDA); elsewhere fall to the host dequant +
   // bf16 Matmul reference. Byte-identical to the old `device == kCUDA` test.
@@ -454,5 +499,15 @@ inline DBuf MatmulNvfp4W4A16D(Dev d, const Tensor& x, const Nvfp4Weight& w,
   return dout;
 }
 
+// Named MXFP4 W4A16 entry point (mirrors MatmulNvfp4W4A16D). y[M,N] = x[M,K] @
+// dequant_mxfp4(w).T for a compressed-tensors MXFP4 weight-only weight (E8M0,
+// group 32, no global). Routes through the SAME shared Marlin/CPU dispatcher via
+// w.is_mxfp4; the assert documents the contract.
+inline DBuf MatmulMxfp4W4A16D(Dev d, const Tensor& x, const Nvfp4Weight& w,
+                              DType out_dtype) {
+  VT_CHECK(w.is_mxfp4, "dense_mxfp4: non-MXFP4 weight routed into MatmulMxfp4W4A16D");
+  return MatmulNvfp4W4A16D(d, x, w, out_dtype);
+}
+
 }  // namespace dense_nvfp4
 }  // namespace vllm
diff --git a/include/vllm/model_executor/models/dense_weight_loaders.h b/include/vllm/model_executor/models/dense_weight_loaders.h
index cd89a54c5..e5040fdf5 100644
--- a/include/vllm/model_executor/models/dense_weight_loaders.h
+++ b/include/vllm/model_executor/models/dense_weight_loaders.h
@@ -380,5 +380,126 @@ inline Nvfp4Weight LoadMergedCtNvfp4W4A16(
   return merged;
 }
 
+// --- compressed-tensors MXFP4 **W4A16** (`mxfp4-pack-quantized`) -------------
+// ADDED (append-only; no existing helper touched) by the QUANT-CT-MXFP4 dense-Qwen
+// bring-up. MXFP4 analog of LoadCtNvfp4W4A16: same [N=out, K=in] raw orientation
+// and merged-shard ownership rule, but the block scale is E8M0/UE8M0 at
+// group_size 32 with NO global scale.
+//
+// ON-DISK LAYOUT (verified on Yi30/Qwen3-8B-MXFP4, the format compressed-tensors
+// emits for `mxfp4-pack-quantized`, group_size 32, num_bits 4, type float):
+//   <proj>.weight_packed  U8  [N, K/2]    two E2M1 nibbles per byte
+//   <proj>.weight_scale   U8  [N, K/32]   one E8M0 (biased exponent) per 32 elems
+// and NO `<proj>.weight_global_scale` / `<proj>.input_global_scale` (MXFP4 has no
+// global, and the FlashInfer/Marlin W4A16 path folds the E8M0 scale directly). On
+// GB10 the runnable oracle path is Marlin W4A16 (FlashInfer cute-dsl mxf4 rejects
+// sm_121), which is exactly what this routes to. The result is a raw fp4-resident
+// Nvfp4Weight with is_mxfp4=true, group_size=32, scale2 unused.
+
+// True when `proj` is stored as a compressed-tensors MXFP4 linear: `.weight_packed`
+// present AND `.weight_scale` is U8 (E8M0). The U8 scale is the discriminator vs
+// NVFP4 (whose weight_scale is F8_E4M3), so this never matches an NVFP4 checkpoint.
+inline bool IsCtMxfp4Projection(
+    const TensorResolver& get,
+    const std::function<bool(const std::string&)>& has, const std::string& proj) {
+  if (!has(proj + ".weight_packed") || !has(proj + ".weight_scale")) return false;
+  return get(proj + ".weight_scale").dtype == "U8";
+}
+
+// One compressed-tensors MXFP4 W4A16 Linear -> raw fp4-resident Nvfp4Weight in the
+// on-disk [N=out, K=in] orientation the Marlin W4A16 GEMM reads directly.
+// `has` is unused (MXFP4 carries no optional global/input-scale tensors to probe)
+// but kept for signature parity with LoadCtNvfp4W4A16 so the merged loader and the
+// per-projection call sites are uniform.
+inline Nvfp4Weight LoadCtMxfp4W4A16(
+    const TensorResolver& get,
+    [[maybe_unused]] const std::function<bool(const std::string&)>& has,
+    const std::string& proj) {
+  const StTensor& packed = get(proj + ".weight_packed");
+  VT_CHECK(packed.dtype == "U8",
+           "dense loader: expected U8 weight_packed for " + proj);
+  VT_CHECK(packed.shape.size() == 2,
+           "dense loader: expected 2-D weight_packed for " + proj);
+  const int64_t out_dim = packed.shape[0];
+  const int64_t in_dim = packed.shape[1] * 2;
+  VT_CHECK(in_dim % 32 == 0,
+           "dense loader: MXFP4 in_dim must be a multiple of 32 for " + proj);
+  const StTensor& ws = get(proj + ".weight_scale");
+  VT_CHECK(ws.dtype == "U8",
+           "dense loader: expected U8 (E8M0) weight_scale for " + proj);
+  VT_CHECK(ws.shape.size() == 2 && ws.shape[0] == out_dim &&
+               ws.shape[1] == in_dim / 32,
+           "dense loader: weight_scale shape must be [N, K/32] for " + proj);
+
+  Nvfp4Weight r;
+  r.n = out_dim;
+  r.k = in_dim;
+  r.group_size = 32;
+  r.is_mxfp4 = true;
+  r.scale2 = 0.0F;  // MXFP4 has no global scale (unused)
+  r.alpha = 0.0F;   // W4A16: no activation quant
+  r.packed = MakeOwned(vt::DType::kI8, {out_dim, in_dim / 2});
+  VT_CHECK(packed.nbytes == r.packed.bytes.size(),
+           "dense loader: packed byte-size mismatch for " + proj);
+  std::memcpy(r.packed.bytes.data(), packed.data, packed.nbytes);
+  MaybeReleaseSourcePages(packed.data, packed.nbytes);
+  r.scale = MakeOwned(vt::DType::kI8, {out_dim, in_dim / 32});
+  VT_CHECK(ws.nbytes == r.scale.bytes.size(),
+           "dense loader: scale byte-size mismatch for " + proj);
+  std::memcpy(r.scale.bytes.data(), ws.data, ws.nbytes);
+  MaybeReleaseSourcePages(ws.data, ws.nbytes);
+  return r;
+}
+
+// Load and concatenate compressed-tensors MXFP4 W4A16 shards `[N_i, K]` along
+// output rows (the MXFP4 analog of LoadMergedCtNvfp4W4A16). Both weight_packed
+// [N_i,K/2] and weight_scale [N_i,K/32] are row-major over N, so both concat by
+// plain row-stack (grouping runs along K, untouched). No global scale to collapse.
+inline Nvfp4Weight LoadMergedCtMxfp4W4A16(
+    const TensorResolver& get, const std::function<bool(const std::string&)>& has,
+    const std::vector<std::string>& projs) {
+  VT_CHECK(!projs.empty(),
+           "dense loader: merged MXFP4 projection requires at least one shard");
+  std::vector<Nvfp4Weight> shards;
+  shards.reserve(projs.size());
+  int64_t in_dim = -1;
+  int64_t out_dim = 0;
+  for (const std::string& proj : projs) {
+    Nvfp4Weight s = LoadCtMxfp4W4A16(get, has, proj);
+    if (in_dim < 0) in_dim = s.k;
+    VT_CHECK(s.k == in_dim,
+             "dense loader: merged MXFP4 shards must share input width");
+    VT_CHECK(out_dim <= std::numeric_limits<int64_t>::max() - s.n,
+             "dense loader: merged MXFP4 output width overflow");
+    out_dim += s.n;
+    shards.push_back(std::move(s));
+  }
+  if (shards.size() == 1) return std::move(shards[0]);
+
+  Nvfp4Weight merged;
+  merged.n = out_dim;
+  merged.k = in_dim;
+  merged.group_size = 32;
+  merged.is_mxfp4 = true;
+  merged.scale2 = 0.0F;
+  merged.alpha = 0.0F;
+  merged.packed = MakeOwned(vt::DType::kI8, {out_dim, in_dim / 2});
+  merged.scale = MakeOwned(vt::DType::kI8, {out_dim, in_dim / 32});
+  size_t p_off = 0;
+  size_t s_off = 0;
+  for (const Nvfp4Weight& s : shards) {
+    std::memcpy(merged.packed.bytes.data() + p_off, s.packed.bytes.data(),
+                s.packed.bytes.size());
+    p_off += s.packed.bytes.size();
+    std::memcpy(merged.scale.bytes.data() + s_off, s.scale.bytes.data(),
+                s.scale.bytes.size());
+    s_off += s.scale.bytes.size();
+  }
+  VT_CHECK(p_off == merged.packed.bytes.size() &&
+               s_off == merged.scale.bytes.size(),
+           "dense loader: merged MXFP4 byte accounting mismatch");
+  return merged;
+}
+
 }  // namespace dense_loaders
 }  // namespace vllm
diff --git a/include/vllm/model_executor/models/qwen3_5_weights.h b/include/vllm/model_executor/models/qwen3_5_weights.h
index 8e34eecd3..8563aea88 100644
--- a/include/vllm/model_executor/models/qwen3_5_weights.h
+++ b/include/vllm/model_executor/models/qwen3_5_weights.h
@@ -117,6 +117,15 @@ struct Nvfp4Weight {
   int64_t k = 0;        // in_features (K % 16 == 0)
   bool Empty() const { return packed.Empty(); }
 
+  // Block-scale FORMAT. Default = NVFP4: group_size 16, fp8-e4m3 `scale`, a
+  // per-tensor `scale2` global. is_mxfp4 selects compressed-tensors MXFP4
+  // (`mxfp4-pack-quantized`): group_size 32, E8M0 (UE8M0) `scale` [N, K/32], NO
+  // global (scale2 unused). Same E2M1 `packed`. Set ONLY by the dense MXFP4
+  // loader; the 27B/35B NVFP4 gate paths never touch it (default false) so their
+  // Marlin repack/GEMM are byte-unchanged.
+  int group_size = 16;
+  bool is_mxfp4 = false;
+
   // TRUE W4A4 fields (27B compressed-tensors NVFP4; notes §7). Populated ONLY on
   // the 27B CT load (LoadCtNvfp4Raw); left 0 for the 35B modelopt W4A16 weights
   // (which have no activation quant) so `IsTrueW4A4()` gates the 27B alone.
diff --git a/include/vt/cuda/marlin_repack.h b/include/vt/cuda/marlin_repack.h
index 04509099a..3a62a214f 100644
--- a/include/vt/cuda/marlin_repack.h
+++ b/include/vt/cuda/marlin_repack.h
@@ -21,6 +21,18 @@ void MarlinRepackExpertWeight(void* stream, int device, uint32_t* out_weight,
 void MarlinProcessExpertScales(void* stream, const uint8_t* scale_nk16, uint8_t* out_scale,
                                int size_k, int size_n, float scale_factor);
 
+// MXFP4 variant: process one expert's E8M0 (UE8M0) block scales
+// [N=out, K/32] uint8 (device, group_size 32) into Marlin's permuted E8M0 scale
+// layout [K/32, N] uint8 (device, caller-allocated). Mirrors vLLM
+// mxfp4_marlin_process_scales (marlin_utils_fp4.py:125-139, bf16 activation):
+// marlin_permute_scales(group_size=32) + the within-4 [0,2,1,3] "fit the fp8
+// dequant layout" reorder, then keep the value as an E8M0 byte (NO scale_factor,
+// NO *128/S0E5M3, NO global — the `.to(bf16).to(e8m0)` round-trip is identity for
+// the exact-pow2 E8M0 values). The Marlin kernel decodes each byte as
+// 2^(byte-127) (marlin_template.h:1497-1510). No combined scale factor, no global.
+void MarlinProcessExpertScalesMxfp4(void* stream, const uint8_t* scale_nk32,
+                                    uint8_t* out_scale, int size_k, int size_n);
+
 // combined_scale_factor across a set of experts' host fp8 scale buffers (all
 // experts of one Marlin GEMM share it: gate+up for w13, down alone for w2).
 float MarlinNvfp4CombinedScaleFactor(const std::vector<const uint8_t*>& host_scale_bufs,
diff --git a/include/vt/ops.h b/include/vt/ops.h
index 3b0b95636..6ac0815f3 100644
--- a/include/vt/ops.h
+++ b/include/vt/ops.h
@@ -728,6 +728,12 @@ struct MoeMarlinArgs {
   int size_n = 0;  // output features
   int size_k = 0;  // input features (contraction; multiple of 16)
   bool mul_topk_weights = false;  // fold topk_weights into the output (down proj)
+  // Block-scale format selector. Default = NVFP4 (fp8-e4m3 scales, group 16,
+  // per-tensor global scale). group_size 32 + mxfp4=true selects the MXFP4 path
+  // (E8M0/UE8M0 scales => s_type kFE8M0fnu, group_blocks 2, NO global scale;
+  // the `global_scale` tensor is ignored). Mirrors vLLM's is_nvfp4 branch.
+  int group_size = 16;
+  bool mxfp4 = false;
 };
 using MoeGroupedGemmNvfp4MarlinFn =
     void (*)(Queue&, Tensor&, const Tensor&, const Tensor&, const Tensor&, const Tensor&, Tensor&,
diff --git a/src/vllm/model_executor/models/qwen3_weights.cpp b/src/vllm/model_executor/models/qwen3_weights.cpp
index b13df8646..9e7d18247 100644
--- a/src/vllm/model_executor/models/qwen3_weights.cpp
+++ b/src/vllm/model_executor/models/qwen3_weights.cpp
@@ -40,11 +40,14 @@
 namespace vllm {
 namespace {
 
+using dense_loaders::IsCtMxfp4Projection;
 using dense_loaders::IsCtNvfp4Projection;
 using dense_loaders::LoadBf16Direct;
 using dense_loaders::LoadBf16Transposed;
+using dense_loaders::LoadCtMxfp4W4A16;
 using dense_loaders::LoadCtNvfp4W4A16;
 using dense_loaders::LoadMergedBf16RawNK;
+using dense_loaders::LoadMergedCtMxfp4W4A16;
 using dense_loaders::LoadMergedCtNvfp4W4A16;
 
 using TensorExists = std::function<bool(const std::string&)>;
@@ -74,7 +77,12 @@ Qwen3DenseLayerWeights LoadQwen3Layer(const TensorResolver& get,
   // `<proj>.weight_packed` instead of `<proj>.weight`. The norms and the embed
   // table stay BF16 either way (they are not Linears, so no config group targets
   // them), and `lm_head` is in the checkpoint's `ignore` list.
-  const bool fp4 = IsCtNvfp4Projection(has, sa + "q_proj");
+  // MXFP4 (`mxfp4-pack-quantized`, U8 E8M0 scale) and NVFP4 (`nvfp4-pack-quantized`,
+  // F8_E4M3 scale) both store `<proj>.weight_packed`; the U8 scale discriminates
+  // MXFP4. Both populate the SAME fp4 fields (Nvfp4Weight, is_mxfp4 set for MXFP4),
+  // so the forward routes on the weight flag with no per-call arch probe.
+  const bool mxfp4 = IsCtMxfp4Projection(get, has, sa + "q_proj");
+  const bool fp4 = mxfp4 || IsCtNvfp4Projection(has, sa + "q_proj");
 
   // QKVParallelLinear: one merged owner in exact [q,k,v] output-row order
   // (packed_modules_mapping qkv_proj<-[q,k,v]_proj), kept raw-NK for MatmulBT.
@@ -82,9 +90,15 @@ Qwen3DenseLayerWeights LoadQwen3Layer(const TensorResolver& get,
   // packed + scale, max-then-reciprocate global scale) — the same single merged
   // parameter vLLM owns.
   if (fp4) {
-    w.attn.qkv_proj_fp4 = LoadMergedCtNvfp4W4A16(
-        get, has, {sa + "q_proj", sa + "k_proj", sa + "v_proj"});
-    w.attn.o_proj_fp4 = LoadCtNvfp4W4A16(get, has, sa + "o_proj");
+    if (mxfp4) {
+      w.attn.qkv_proj_fp4 = LoadMergedCtMxfp4W4A16(
+          get, has, {sa + "q_proj", sa + "k_proj", sa + "v_proj"});
+      w.attn.o_proj_fp4 = LoadCtMxfp4W4A16(get, has, sa + "o_proj");
+    } else {
+      w.attn.qkv_proj_fp4 = LoadMergedCtNvfp4W4A16(
+          get, has, {sa + "q_proj", sa + "k_proj", sa + "v_proj"});
+      w.attn.o_proj_fp4 = LoadCtNvfp4W4A16(get, has, sa + "o_proj");
+    }
   } else {
     w.attn.qkv_proj = LoadMergedBf16RawNK(
         get, {sa + "q_proj.weight", sa + "k_proj.weight", sa + "v_proj.weight"});
@@ -104,9 +118,15 @@ Qwen3DenseLayerWeights LoadQwen3Layer(const TensorResolver& get,
   // vLLM's single merged Marlin operand (size_n=2I) at repack time, or runs the
   // split two-GEMM A/B — both from the same loaded bytes.
   if (fp4) {
-    w.mlp.gate_proj_fp4 = LoadCtNvfp4W4A16(get, has, mlp + "gate_proj");
-    w.mlp.up_proj_fp4 = LoadCtNvfp4W4A16(get, has, mlp + "up_proj");
-    w.mlp.down_proj_fp4 = LoadCtNvfp4W4A16(get, has, mlp + "down_proj");
+    if (mxfp4) {
+      w.mlp.gate_proj_fp4 = LoadCtMxfp4W4A16(get, has, mlp + "gate_proj");
+      w.mlp.up_proj_fp4 = LoadCtMxfp4W4A16(get, has, mlp + "up_proj");
+      w.mlp.down_proj_fp4 = LoadCtMxfp4W4A16(get, has, mlp + "down_proj");
+    } else {
+      w.mlp.gate_proj_fp4 = LoadCtNvfp4W4A16(get, has, mlp + "gate_proj");
+      w.mlp.up_proj_fp4 = LoadCtNvfp4W4A16(get, has, mlp + "up_proj");
+      w.mlp.down_proj_fp4 = LoadCtNvfp4W4A16(get, has, mlp + "down_proj");
+    }
   } else {
     w.mlp.gate_up_proj = LoadMergedBf16RawNK(
         get, {mlp + "gate_proj.weight", mlp + "up_proj.weight"});
diff --git a/src/vt/cuda/cuda_marlin_repack.cu b/src/vt/cuda/cuda_marlin_repack.cu
index c1cb87fe6..474986b9f 100644
--- a/src/vt/cuda/cuda_marlin_repack.cu
+++ b/src/vt/cuda/cuda_marlin_repack.cu
@@ -110,6 +110,30 @@ __global__ void ProcessScalesKernel(uint8_t* __restrict__ out,
   out[static_cast<size_t>(r) * SN + c] = obyte;
 }
 
+// MXFP4 E8M0 scale processing. Same permute chain as ProcessScalesKernel
+// (marlin_permute_scales 64-wide scale_perm + within-4 [0,2,1,3]), but over the
+// group-32 scale grid [SK32, SN] and the value is the SOURCE E8M0 byte passed
+// through unchanged (no fp8 decode, no *scale_factor*128, no S0E5M3 encode, no
+// <2 clamp). scalebuf is our [N, K/32] E8M0 block-scale (row-major over N).
+__global__ void ProcessScalesMxfp4Kernel(uint8_t* __restrict__ out,
+                                         const uint8_t* __restrict__ scalebuf,
+                                         int size_k, int size_n) {
+  const int SN = size_n;
+  const int SK32 = size_k / 32;
+  const int r = blockIdx.y * blockDim.y + threadIdx.y;
+  const int c = blockIdx.x * blockDim.x + threadIdx.x;
+  if (r >= SK32 || c >= SN) return;
+
+  const int i = r * SN + c;  // flat index into the [SK32, SN] processed tensor
+  const int perm4[4] = {0, 2, 1, 3};
+  const int mi = (i / 4) * 4 + perm4[i % 4];             // undo within-4 reorder
+  const int j = (mi / 64) * 64 + kScalePerm[mi % 64];    // undo permute_scales
+  const int kr = j / SN;
+  const int n = j % SN;
+  // Source E8M0 byte at scale[n][kr] = scalebuf[n*SK32 + kr]; passthrough.
+  out[static_cast<size_t>(r) * SN + c] = scalebuf[static_cast<size_t>(n) * SK32 + kr];
+}
+
 // fp8-e4m3 byte -> float (host, for the combined_scale_factor max reduce).
 float DecodeFp8E4m3(uint8_t b) {
   const int sign = (b >> 7) & 1;
@@ -175,6 +199,23 @@ void MarlinProcessExpertScales(void* stream, const uint8_t* scale_nk16, uint8_t*
   RCheck(cudaGetLastError(), "process scales launch");
 }
 
+void MarlinProcessExpertScalesMxfp4(void* stream, const uint8_t* scale_nk32,
+                                    uint8_t* out_scale, int size_k, int size_n) {
+  cudaStream_t s = static_cast<cudaStream_t>(stream);
+  static bool perm_uploaded = false;
+  if (!perm_uploaded) {
+    RCheck(cudaMemcpyToSymbol(kScalePerm, kScalePermHost, sizeof(kScalePermHost)),
+           "upload scale_perm (mxfp4)");
+    perm_uploaded = true;
+  }
+  const int SN = size_n;
+  const int SK32 = size_k / 32;
+  dim3 tb(32, 8);
+  dim3 tg((SN + tb.x - 1) / tb.x, (SK32 + tb.y - 1) / tb.y);
+  ProcessScalesMxfp4Kernel<<<tg, tb, 0, s>>>(out_scale, scale_nk32, size_k, size_n);
+  RCheck(cudaGetLastError(), "process scales mxfp4 launch");
+}
+
 float MarlinNvfp4CombinedScaleFactor(const std::vector<const uint8_t*>& host_scale_bufs,
                                      const std::vector<size_t>& lens) {
   // _nvfp4_compute_scale_factor: ws = max(scale)*2^7; if ws < 448*2^7:
diff --git a/src/vt/cuda/cuda_moe_marlin.cu b/src/vt/cuda/cuda_moe_marlin.cu
index a2a26f67c..3a89b6ef9 100644
--- a/src/vt/cuda/cuda_moe_marlin.cu
+++ b/src/vt/cuda/cuda_moe_marlin.cu
@@ -111,11 +111,13 @@ void MoeGroupedGemmNvfp4MarlinKernelCuda(Queue& q, Tensor& c, const Tensor& a,
   cudaStream_t s = AsStream(q);
   const int dev = q.device.index;
 
-  // NVFP4 W4A16, bf16 activation/output (marlin_moe_wna16 generate_kernels.py:94).
+  // NVFP4 W4A16, bf16 activation/output (marlin_moe_wna16 generate_kernels.py:94),
+  // OR MXFP4 W4A16 when args.mxfp4 (E8M0 scales => s_type kFE8M0fnu, group_size 32
+  // => group_blocks 2, NO global scale). Mirrors vLLM's is_nvfp4 branch.
   const vllm::ScalarType a_type = vllm::kBFloat16;
   const vllm::ScalarType b_type = vllm::kFE2M1f;
   const vllm::ScalarType c_type = vllm::kBFloat16;
-  const vllm::ScalarType s_type = vllm::kFE4M3fn;
+  const vllm::ScalarType s_type = args.mxfp4 ? vllm::kFE8M0fnu : vllm::kFE4M3fn;
 
   const int num_experts = static_cast<int>(b_q_weight.shape[0]);
   const int size_m = args.size_m;
@@ -123,8 +125,11 @@ void MoeGroupedGemmNvfp4MarlinKernelCuda(Queue& q, Tensor& c, const Tensor& a,
   const int size_k = args.size_k;
   const int moe_block_size = args.moe_block_size;
   const int top_k = args.top_k;
-  const int group_size = 16;                 // group_blocks == 1
+  const int group_size = args.group_size;    // 16 (nvfp4) or 32 (mxfp4)
   const int num_groups = size_k / group_size;
+  // MXFP4 has NO global scale — the kernel only reads global_scale_ptr under
+  // (b_type==kFE2M1f && s_type==kFE4M3fn), so pass nullptr on the mxfp4 path.
+  void* global_scale_ptr = args.mxfp4 ? nullptr : global_scale.data;
 
   int sms = -1;
   Check(cudaDeviceGetAttribute(&sms, cudaDevAttrMultiProcessorCount, dev),
@@ -153,7 +158,7 @@ void MoeGroupedGemmNvfp4MarlinKernelCuda(Queue& q, Tensor& c, const Tensor& a,
 
   MARLIN_NAMESPACE_NAME::marlin_mm(
       a.data, b_q_weight.data, c.data, c_tmp, /*b_bias=*/nullptr, /*a_s=*/nullptr,
-      b_scales.data, global_scale.data, /*zp=*/nullptr, /*g_idx=*/nullptr, /*perm=*/nullptr,
+      b_scales.data, global_scale_ptr, /*zp=*/nullptr, /*g_idx=*/nullptr, /*perm=*/nullptr,
       /*a_tmp=*/nullptr, sorted_token_ids.data, expert_ids.data, num_tokens_past_padded.data,
       topk_weights.data, moe_block_size, num_experts, top_k, args.mul_topk_weights, size_m,
       size_n, size_k, workspace.data, a_type, b_type, c_type, s_type, /*has_bias=*/false,
diff --git a/src/vt/cuda/marlin/libtorch_stable/moe/marlin_moe_wna16/generate_kernels.py b/src/vt/cuda/marlin/libtorch_stable/moe/marlin_moe_wna16/generate_kernels.py
index 6b3b2530f..8182e90f9 100644
--- a/src/vt/cuda/marlin/libtorch_stable/moe/marlin_moe_wna16/generate_kernels.py
+++ b/src/vt/cuda/marlin/libtorch_stable/moe/marlin_moe_wna16/generate_kernels.py
@@ -72,6 +72,19 @@
         "thread_m_blocks": THREAD_M_BLOCKS,
         "group_blocks": [1],
     },
+    # MXFP4 (compressed-tensors mxfp4-pack-quantized, bf16 activation) — the
+    # QUANT-CT-MXFP4 dense-Qwen W4A16 path. Same E2M1 weight, but E8M0 (UE8M0)
+    # block scales at group_size 32 => group_blocks == 2, and NO global scale.
+    # Mirrors vLLM marlin_utils_fp4.prepare_fp4_layer_for_marlin (is_nvfp4=False:
+    # group_size 32, mxfp4_marlin_process_scales, weight_global_scale=None) @ pin.
+    {
+        "a_type": ["kBFloat16"],
+        "b_type": "kFE2M1f",
+        "s_type": "kFE8M0fnu",
+        "thread_configs": THREAD_CONFIGS,
+        "thread_m_blocks": THREAD_M_BLOCKS,
+        "group_blocks": [2],
+    },
 ]
 
 
diff --git a/src/vt/cuda/marlin/libtorch_stable/moe/marlin_moe_wna16/kernel_selector.h b/src/vt/cuda/marlin/libtorch_stable/moe/marlin_moe_wna16/kernel_selector.h
index fd6905de1..0f8c19390 100644
--- a/src/vt/cuda/marlin/libtorch_stable/moe/marlin_moe_wna16/kernel_selector.h
+++ b/src/vt/cuda/marlin/libtorch_stable/moe/marlin_moe_wna16/kernel_selector.h
@@ -30,3 +30,33 @@ else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm:
   kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE4M3fn.id(), 128, 4, 8, 4, false, 4, 1, false>;
 else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE4M3fn && threads == 128 && thread_m_blocks == 4 && thread_n_blocks == 4 && thread_k_blocks == 8 && m_block_size_8 == false && stages == 4 && group_blocks == 1 && is_zp_float == false)
   kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE4M3fn.id(), 128, 4, 4, 8, false, 4, 1, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 256 && thread_m_blocks == 1 && thread_n_blocks == 8 && thread_k_blocks == 8 && m_block_size_8 == true && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 256, 1, 8, 8, true, 4, 2, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 128 && thread_m_blocks == 1 && thread_n_blocks == 8 && thread_k_blocks == 4 && m_block_size_8 == true && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 1, 8, 4, true, 4, 2, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 128 && thread_m_blocks == 1 && thread_n_blocks == 4 && thread_k_blocks == 8 && m_block_size_8 == true && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 1, 4, 8, true, 4, 2, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 256 && thread_m_blocks == 1 && thread_n_blocks == 8 && thread_k_blocks == 8 && m_block_size_8 == false && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 256, 1, 8, 8, false, 4, 2, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 128 && thread_m_blocks == 1 && thread_n_blocks == 8 && thread_k_blocks == 4 && m_block_size_8 == false && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 1, 8, 4, false, 4, 2, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 128 && thread_m_blocks == 1 && thread_n_blocks == 4 && thread_k_blocks == 8 && m_block_size_8 == false && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 1, 4, 8, false, 4, 2, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 256 && thread_m_blocks == 2 && thread_n_blocks == 16 && thread_k_blocks == 4 && m_block_size_8 == false && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 256, 2, 16, 4, false, 4, 2, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 128 && thread_m_blocks == 2 && thread_n_blocks == 8 && thread_k_blocks == 4 && m_block_size_8 == false && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 2, 8, 4, false, 4, 2, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 128 && thread_m_blocks == 2 && thread_n_blocks == 4 && thread_k_blocks == 8 && m_block_size_8 == false && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 2, 4, 8, false, 4, 2, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 256 && thread_m_blocks == 3 && thread_n_blocks == 16 && thread_k_blocks == 4 && m_block_size_8 == false && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 256, 3, 16, 4, false, 4, 2, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 128 && thread_m_blocks == 3 && thread_n_blocks == 8 && thread_k_blocks == 4 && m_block_size_8 == false && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 3, 8, 4, false, 4, 2, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 128 && thread_m_blocks == 3 && thread_n_blocks == 4 && thread_k_blocks == 8 && m_block_size_8 == false && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 3, 4, 8, false, 4, 2, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 256 && thread_m_blocks == 4 && thread_n_blocks == 16 && thread_k_blocks == 4 && m_block_size_8 == false && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 256, 4, 16, 4, false, 4, 2, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 128 && thread_m_blocks == 4 && thread_n_blocks == 8 && thread_k_blocks == 4 && m_block_size_8 == false && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 4, 8, 4, false, 4, 2, false>;
+else if (a_type == vllm::kBFloat16 && b_type == vllm::kFE2M1f && c_type == vllm::kBFloat16 && s_type == vllm::kFE8M0fnu && threads == 128 && thread_m_blocks == 4 && thread_n_blocks == 4 && thread_k_blocks == 8 && m_block_size_8 == false && stages == 4 && group_blocks == 2 && is_zp_float == false)
+  kernel = Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 4, 4, 8, false, 4, 2, false>;
diff --git a/src/vt/cuda/marlin/libtorch_stable/moe/marlin_moe_wna16/sm80_kernel_bfloat16_fe2m1f_bfloat16.cu b/src/vt/cuda/marlin/libtorch_stable/moe/marlin_moe_wna16/sm80_kernel_bfloat16_fe2m1f_bfloat16.cu
index 63e98e79c..8688df862 100644
--- a/src/vt/cuda/marlin/libtorch_stable/moe/marlin_moe_wna16/sm80_kernel_bfloat16_fe2m1f_bfloat16.cu
+++ b/src/vt/cuda/marlin/libtorch_stable/moe/marlin_moe_wna16/sm80_kernel_bfloat16_fe2m1f_bfloat16.cu
@@ -37,4 +37,34 @@ template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::
 
 template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE4M3fn.id(), 128, 4, 4, 8, false, 4, 1, false>( MARLIN_KERNEL_PARAMS );
 
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 256, 1, 8, 8, true, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 1, 8, 4, true, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 1, 4, 8, true, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 256, 1, 8, 8, false, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 1, 8, 4, false, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 1, 4, 8, false, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 256, 2, 16, 4, false, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 2, 8, 4, false, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 2, 4, 8, false, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 256, 3, 16, 4, false, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 3, 8, 4, false, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 3, 4, 8, false, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 256, 4, 16, 4, false, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 4, 8, 4, false, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
+template __global__ void Marlin<vllm::kBFloat16.id(), vllm::kFE2M1f.id(), vllm::kBFloat16.id(), vllm::kFE8M0fnu.id(), 128, 4, 4, 8, false, 4, 2, false>( MARLIN_KERNEL_PARAMS );
+
 }
diff --git a/tests/vllm/model_executor/layers/test_linear_method.cpp b/tests/vllm/model_executor/layers/test_linear_method.cpp
index 95ae90ef5..9ccf68fc3 100644
--- a/tests/vllm/model_executor/layers/test_linear_method.cpp
+++ b/tests/vllm/model_executor/layers/test_linear_method.cpp
@@ -16,9 +16,13 @@
 #include <vector>
 
 #include "vllm/model_executor/layers/quantization/compressed_tensors/schemes/nvfp4.h"
+#include "vllm/model_executor/model_loader/mxfp4_dequant.h"
 #include "vt/backend.h"
 #include "vt/dtype.h"
 
+#include <cmath>
+#include <random>
+
 namespace {
 
 using vllm::Nvfp4Weight;
@@ -67,8 +71,116 @@ Nvfp4Weight MakeNvfp4W4A16(int64_t N, int64_t K) {
   return w;
 }
 
+// A random MXFP4 W4A16 weight: E2M1 packed [N,K/2] + E8M0 scale [N,K/32], group
+// 32, no global, is_mxfp4=true — so the factory + Apply route the MXFP4 keep-quant
+// path (Marlin on GPU via BuildMarlinDenseResident).
+Nvfp4Weight MakeMxfp4W4A16(int64_t N, int64_t K, uint32_t seed) {
+  Nvfp4Weight w;
+  w.n = N;
+  w.k = K;
+  w.group_size = 32;
+  w.is_mxfp4 = true;
+  w.scale2 = 0.0f;
+  w.packed.dtype = DType::kI8;
+  w.packed.rank = 2;
+  w.packed.shape[0] = N;
+  w.packed.shape[1] = K / 2;
+  w.packed.bytes.resize(static_cast<size_t>(N) * (K / 2));
+  w.scale.dtype = DType::kI8;
+  w.scale.rank = 2;
+  w.scale.shape[0] = N;
+  w.scale.shape[1] = K / 32;
+  w.scale.bytes.resize(static_cast<size_t>(N) * (K / 32));
+  std::mt19937 rng(seed);
+  for (auto& b : w.packed.bytes) b = static_cast<uint8_t>(rng() & 0xFFu);
+  for (auto& s : w.scale.bytes) s = static_cast<uint8_t>(118u + (rng() % 15u));
+  return w;
+}
+
 }  // namespace
 
+#ifdef VT_MARLIN_NVFP4
+// The model-facing MXFP4 path END-TO-END: MakeLinearMethod(bf16-empty, mxfp4) ->
+// Apply -> MatmulNvfp4W4A16D -> (GPU) MatmulNvfp4MarlinD -> BuildMarlinDenseResident
+// -> MoeGroupedGemmNvfp4Marlin. This is the ONE link the op-level unit gate does NOT
+// cover (it feeds MANUALLY-built residents), so it isolates a resident-builder bug
+// from the kernel. Reference = the INDEPENDENT CPU dequant (DequantMxfp4ToF32 + f32
+// matmul). Real Qwen3-8B projection shapes; M=1 (decode) AND M=8 (prefill).
+TEST_CASE("linear_method: MXFP4 W4A16 Apply (Marlin BuildMarlinDenseResident) == CPU dequant ref") {
+  vt::Backend* gpu = nullptr;
+  try {
+    gpu = &vt::GetBackend(vt::DeviceType::kCUDA);
+  } catch (...) {
+    MESSAGE("SKIP: no CUDA backend");
+    return;
+  }
+  // Persist weights in a vector so each shape has a DISTINCT, stable address —
+  // the resident cache (MarlinDenseResidentFor) is keyed by weight pointer, and a
+  // loop-local reused stack slot would alias residents across shapes (a test
+  // artifact, not a model bug: the model's weights are distinct persistent objects).
+  const std::vector<std::pair<int64_t, int64_t>> shapes{{4096, 4096}, {12288, 4096}};
+  std::vector<Nvfp4Weight> weights;
+  for (auto KN : shapes) weights.push_back(MakeMxfp4W4A16(KN.second, KN.first, 2024));
+  for (size_t si = 0; si < shapes.size(); ++si) {
+    const int64_t K = shapes[si].first, N = shapes[si].second;
+    CAPTURE(K);
+    CAPTURE(N);
+    Nvfp4Weight& w = weights[si];
+    OwnedTensor bf16_empty;  // Empty() => factory selects the fp4 method
+
+    std::vector<float> w_f32(static_cast<size_t>(N * K));
+    vllm::DequantMxfp4ToF32(reinterpret_cast<const uint8_t*>(w.packed.bytes.data()),
+                            reinterpret_cast<const uint8_t*>(w.scale.bytes.data()), N, K,
+                            w_f32.data());
+
+    for (int64_t M : {int64_t{1}, int64_t{8}}) {
+      CAPTURE(M);
+      vt::Queue q = gpu->CreateQueue();
+      vllm::dense_attn::Dev d{*gpu, q};
+
+      std::vector<uint16_t> act_bf16(static_cast<size_t>(M * K));
+      std::mt19937 rng(7 + static_cast<uint32_t>(M));
+      std::uniform_real_distribution<float> dist(-1.0f, 1.0f);
+      std::vector<float> act_r(static_cast<size_t>(M * K));
+      for (size_t i = 0; i < act_bf16.size(); ++i) {
+        act_bf16[i] = vt::F32ToBF16(dist(rng));
+        act_r[i] = vt::BF16ToF32(act_bf16[i]);
+      }
+      std::vector<float> ref(static_cast<size_t>(M * N), 0.0f);
+      for (int64_t m = 0; m < M; ++m)
+        for (int64_t n = 0; n < N; ++n) {
+          float acc = 0.0f;
+          for (int64_t k = 0; k < K; ++k)
+            acc += act_r[static_cast<size_t>(m * K + k)] * w_f32[static_cast<size_t>(n * K + k)];
+          ref[static_cast<size_t>(m * N + n)] = acc;
+        }
+
+      vllm::dense_attn::DBuf x(d, DType::kBF16, {M, K}, act_bf16.data());
+      auto method = layers::MakeLinearMethod(bf16_empty, w);
+      vllm::dense_attn::DBuf out = method->Apply(d, x.t(), DType::kBF16);
+      std::vector<uint16_t> got_bf16(static_cast<size_t>(M * N));
+      gpu->Copy(q, got_bf16.data(), out.t().data,
+                got_bf16.size() * sizeof(uint16_t));
+      gpu->Synchronize(q);
+      double max_rel = 0.0, max_abs = 0.0;
+      size_t bad = 0;
+      for (size_t i = 0; i < got_bf16.size(); ++i) {
+        const float g = vt::BF16ToF32(got_bf16[i]);
+        const float a = std::fabs(g - ref[i]);
+        const float tol = 2e-2f + 2e-2f * std::fabs(ref[i]);
+        if (a > tol) ++bad;
+        max_abs = std::max(max_abs, static_cast<double>(a));
+        max_rel = std::max(max_rel, static_cast<double>(a / (std::fabs(ref[i]) + 1e-6f)));
+      }
+      MESSAGE("MXFP4 Apply K=" << K << " N=" << N << " M=" << M
+              << " bad=" << bad << " max_abs=" << max_abs << " max_rel=" << max_rel);
+      CHECK(bad == 0);
+      gpu->DestroyQueue(q);
+    }
+  }
+}
+#endif  // VT_MARLIN_NVFP4
+
 TEST_CASE("linear_method: factory selects bf16 vs nvfp4-w4a16 by weight presence") {
   OwnedTensor bf16 = MakeBf16({4, 16}, 1);
   Nvfp4Weight empty_fp4;      // Empty() == true
diff --git a/tests/vllm/models/test_qwen3_forward.cpp b/tests/vllm/models/test_qwen3_forward.cpp
index 6a959385f..11b547ee5 100644
--- a/tests/vllm/models/test_qwen3_forward.cpp
+++ b/tests/vllm/models/test_qwen3_forward.cpp
@@ -26,6 +26,7 @@
 #include <string>
 #include <system_error>
 
+#include "vllm/model_executor/model_loader/mxfp4_dequant.h"
 #include "vllm/model_executor/model_loader/nvfp4_dequant.h"
 #include "vllm/model_executor/model_loader/safetensors_reader.h"
 #include "vllm/model_executor/models/dense_nvfp4_gemm.h"
@@ -280,6 +281,41 @@ vllm::OwnedTensor DequantToBf16RawNK(const vllm::Nvfp4Weight& w) {
   return o;
 }
 
+// MXFP4 analog of MakeNvfp4: E2M1 packed [n,k/2] + E8M0 scale [n,k/32] (group 32,
+// no global). is_mxfp4=true so the forward routes it through the MXFP4 keep-quant
+// path (Marlin on GPU, dequant fallback on CPU).
+vllm::Nvfp4Weight MakeMxfp4(int64_t n, int64_t k, uint32_t seed) {
+  REQUIRE(k % 32 == 0);
+  vllm::Nvfp4Weight w;
+  w.n = n;
+  w.k = k;
+  w.group_size = 32;
+  w.is_mxfp4 = true;
+  w.scale2 = 0.0f;  // no global
+  w.alpha = 0.0f;   // W4A16
+  w.packed = vllm::dense_loaders::MakeOwned(DType::kI8, {n, k / 2});
+  w.scale = vllm::dense_loaders::MakeOwned(DType::kI8, {n, k / 32});
+  std::mt19937 rng(seed);
+  auto* p = reinterpret_cast<uint8_t*>(w.packed.bytes.data());
+  for (size_t i = 0; i < w.packed.bytes.size(); ++i)
+    p[i] = static_cast<uint8_t>(rng() & 0xFFu);
+  // E8M0 bytes near 127 => scales ~2^-9..2^3 (finite, bf16 normal range).
+  auto* s = reinterpret_cast<uint8_t*>(w.scale.bytes.data());
+  for (size_t i = 0; i < w.scale.bytes.size(); ++i)
+    s[i] = static_cast<uint8_t>(118u + (rng() % 15u));
+  return w;
+}
+
+vllm::OwnedTensor DequantMxfp4ToBf16RawNK(const vllm::Nvfp4Weight& w) {
+  vllm::OwnedTensor o = vllm::dense_loaders::MakeOwned(DType::kBF16, {w.n, w.k});
+  o.nk = true;
+  vllm::DequantMxfp4ToBf16(
+      reinterpret_cast<const uint8_t*>(w.packed.bytes.data()),
+      reinterpret_cast<const uint8_t*>(w.scale.bytes.data()), w.n, w.k,
+      reinterpret_cast<uint16_t*>(o.bytes.data()));
+  return o;
+}
+
 // Merge two raw-NK bf16 [N,K] tensors by output-row concat (the BF16 arm's
 // gate_up layout, which the fp4 arm keeps as two separate operands).
 vllm::OwnedTensor ConcatRawNK(const vllm::OwnedTensor& a,
@@ -375,6 +411,75 @@ TEST_CASE("qwen3 dense forward: NVFP4 W4A16 == BF16-on-dequantized (CPU syntheti
   CHECK(std::memcmp(a.data(), a2.data(), a.size() * sizeof(float)) == 0);
 }
 
+// MXFP4 analog: the FULL model forward with MXFP4 W4A16 weights must equal the
+// forward on the SAME weights dequantized to BF16 — so any difference is a MXFP4
+// dispatch/layout/merge bug in the model integration, NOT a quantization artifact.
+// On GPU this exercises the Marlin mxf4 path THROUGH the model (merged qkv, split
+// gate_up, o/down) — the exact path the e2e uses, which the op-level unit gate does
+// not (it bypasses the loader/merge/model). RED-first for the e2e residual.
+TEST_CASE("qwen3 dense forward: MXFP4 W4A16 == BF16-on-dequantized (synthetic)") {
+  setenv("VT_FUSED_CHAIN_ADOPT", "1", 1);
+  const HfConfig c = TinyConfig();
+  const int64_t H = c.hidden_size, Hq = c.num_attention_heads;
+  const int64_t Hkv = c.num_key_value_heads, Dh = c.head_dim;
+  const int64_t I = c.intermediate_size;
+  const int64_t qdim = Hq * Dh, kdim = Hkv * Dh;
+
+  Qwen3DenseWeights fp4 = TinyWeights(c);
+  Qwen3DenseWeights deq = TinyWeights(c);
+  uint32_t seed = 9000;
+  for (int64_t l = 0; l < c.num_hidden_layers; ++l) {
+    auto& fl = fp4.layers[static_cast<size_t>(l)];
+    auto& dl = deq.layers[static_cast<size_t>(l)];
+
+    vllm::Nvfp4Weight qkv = MakeMxfp4(qdim + 2 * kdim, H, seed++);
+    vllm::Nvfp4Weight o = MakeMxfp4(H, qdim, seed++);
+    vllm::Nvfp4Weight g = MakeMxfp4(I, H, seed++);
+    vllm::Nvfp4Weight u = MakeMxfp4(I, H, seed++);
+    vllm::Nvfp4Weight dn = MakeMxfp4(H, I, seed++);
+
+    dl.attn.qkv_proj = DequantMxfp4ToBf16RawNK(qkv);
+    dl.attn.o_proj = DequantMxfp4ToBf16RawNK(o);
+    dl.mlp.gate_up_proj =
+        ConcatRawNK(DequantMxfp4ToBf16RawNK(g), DequantMxfp4ToBf16RawNK(u));
+    dl.mlp.down_proj = DequantMxfp4ToBf16RawNK(dn);
+
+    fl.attn.qkv_proj = vllm::OwnedTensor{};
+    fl.attn.o_proj = vllm::OwnedTensor{};
+    fl.mlp.gate_up_proj = vllm::OwnedTensor{};
+    fl.mlp.down_proj = vllm::OwnedTensor{};
+    fl.attn.qkv_proj_fp4 = std::move(qkv);
+    fl.attn.o_proj_fp4 = std::move(o);
+    fl.mlp.gate_proj_fp4 = std::move(g);
+    fl.mlp.up_proj_fp4 = std::move(u);
+    fl.mlp.down_proj_fp4 = std::move(dn);
+
+    REQUIRE(fl.attn.IsNvfp4());  // emptiness-based dispatch; true for mxfp4 too
+    REQUIRE(fl.mlp.IsNvfp4());
+    REQUIRE_FALSE(dl.attn.IsNvfp4());
+    REQUIRE_FALSE(dl.mlp.IsNvfp4());
+  }
+
+  vllm::dense_nvfp4::ResetW4A16Stats();
+  const std::vector<float> a = RunForward(c, fp4);
+  const vllm::dense_nvfp4::Nvfp4W4A16Stats st = vllm::dense_nvfp4::GetW4A16Stats();
+  REQUIRE(a.size() == static_cast<size_t>(5 * c.vocab_size));
+  for (float x : a) REQUIRE(std::isfinite(x));
+  MESSAGE("MXFP4 W4A16 counters: fallback_gemms=" << st.fallback_gemms
+          << " marlin_gemms=" << st.marlin_gemms);
+
+  const std::vector<float> b = RunForward(c, deq);
+  REQUIRE(b.size() == a.size());
+  double max_abs = 0.0;
+  for (size_t i = 0; i < a.size(); ++i)
+    max_abs = std::max(max_abs, static_cast<double>(std::fabs(a[i] - b[i])));
+  MESSAGE("MXFP4 W4A16 vs BF16-on-dequantized: max |dlogit| = " << max_abs);
+  CHECK(max_abs < 1e-2);
+
+  const std::vector<float> a2 = RunForward(c, fp4);
+  CHECK(std::memcmp(a.data(), a2.data(), a.size() * sizeof(float)) == 0);
+}
+
 namespace {
 namespace fs = std::filesystem;
 std::string FindQwen3Snap() {
diff --git a/tests/vt/test_ops_moe_grouped.cpp b/tests/vt/test_ops_moe_grouped.cpp
index 54d4d5abf..e4a2e65e5 100644
--- a/tests/vt/test_ops_moe_grouped.cpp
+++ b/tests/vt/test_ops_moe_grouped.cpp
@@ -16,6 +16,7 @@
 #include <stdexcept>
 #include <vector>
 
+#include "vllm/model_executor/model_loader/mxfp4_dequant.h"
 #include "vllm/model_executor/model_loader/nvfp4_dequant.h"
 #include "vt/cuda/moe_decode_ref.h"
 #include "vt/backend.h"
@@ -140,6 +141,25 @@ Nvfp4Weight MakeNvfp4Weight(int64_t n, int64_t k, uint32_t seed) {
   return w;
 }
 
+// MXFP4 weight: E2M1 packed [N,K/2] + E8M0 (UE8M0) block scale [N,K/32], no global.
+struct Mxfp4Weight {
+  std::vector<uint8_t> packed;  // [N, K/2]
+  std::vector<uint8_t> scale;   // [N, K/32]  E8M0 biased exponent (2^(byte-127))
+};
+
+Mxfp4Weight MakeMxfp4Weight(int64_t n, int64_t k, uint32_t seed) {
+  Mxfp4Weight w;
+  w.packed.resize(static_cast<size_t>(n * (k / 2)));
+  w.scale.resize(static_cast<size_t>(n * (k / 32)));
+  std::mt19937 rng(seed);
+  std::uniform_int_distribution<int> byte_dist(0, 255);
+  // E8M0 bytes near 127 => scales ~2^-9..2^3 (finite, in bf16 normal range).
+  std::uniform_int_distribution<int> e8m0_dist(118, 132);
+  for (auto& b : w.packed) b = static_cast<uint8_t>(byte_dist(rng));
+  for (auto& s : w.scale) s = static_cast<uint8_t>(e8m0_dist(rng));
+  return w;
+}
+
 std::vector<float> RandomF32(size_t numel, uint32_t seed) {
   std::mt19937 rng(seed);
   std::uniform_real_distribution<float> dist(-2.0f, 2.0f);
@@ -737,6 +757,127 @@ TEST_CASE("CUDA marlin fused w13 (size_n=2N) is bit-exact vs split gate/up GEMMs
 // independently), so sorted_ids is compared as a per-expert multiset, not
 // position-for-position. Adversarial inputs: empty experts, one hot expert,
 // uneven loads, ties on block boundaries, and M in {1,8,16}.
+// MXFP4 W4A16 Marlin GEMM (group_blocks=2, E8M0 scales, no global) must match the
+// INDEPENDENT CPU dequant reference (DequantMxfp4ToF32 + f32 matmul) — a different
+// code path than the Marlin repack+kernel dequant, so this is a real cross-check
+// (not a shared-helper tautology). Single expert, all tokens -> expert 0 (the dense
+// MatmulMxfp4W4A16D routing). RED-first coverage of the M=1 DECODE path AND M=8.
+TEST_CASE("CUDA marlin MXFP4 W4A16 (group_blocks=2, E8M0) matches CPU dequant reference") {
+  if (!HasCuda()) {
+    MESSAGE("no CUDA backend registered; skipping");
+    return;
+  }
+  Backend& gpu = vt::GetBackend(DeviceType::kCUDA);
+
+  // K%128==0 (Marlin tile) and %32 (mxfp4 group); N%64==0. Cover the small case
+  // AND the real Qwen3-8B projection shapes (o/qkv K=4096; down K=12288; gate/up
+  // N=12288) so the group_blocks=2 kernel is exercised at model scale, not just a
+  // tiny tile — a large-shape-only bug would pass the tiny case and corrupt e2e.
+  for (auto KN : std::vector<std::pair<int64_t, int64_t>>{
+           {256, 128}, {4096, 4096}, {4096, 12288}, {12288, 4096}}) {
+  const int64_t K = KN.first, N = KN.second;
+  CAPTURE(K);
+  CAPTURE(N);
+  Mxfp4Weight w = MakeMxfp4Weight(N, K, 4321);
+
+  // CPU reference weight [N,K] f32 via the independent dequant helper.
+  std::vector<float> w_f32(static_cast<size_t>(N * K));
+  vllm::DequantMxfp4ToF32(w.packed.data(), w.scale.data(), N, K, w_f32.data());
+
+  for (int64_t M : {int64_t{1}, int64_t{8}}) {
+    CAPTURE(M);
+    QueueGuard gq(gpu);
+    void* stream = gq.q.handle;
+    const int dev = gq.q.device.index;
+    const int64_t top_k = 1, E = 1, P = M * top_k;
+
+    const auto act_f = RandomF32(static_cast<size_t>(M * K), 7000 + static_cast<uint32_t>(M));
+    const auto act_bf16 = ToBf16(act_f);
+    std::vector<float> act_r(act_f.size());
+    for (size_t i = 0; i < act_r.size(); ++i) act_r[i] = vt::BF16ToF32(act_bf16[i]);
+
+    // Reference: out[m,n] = sum_k act_r[m,k] * w_f32[n,k] (f32 accum).
+    std::vector<float> ref(static_cast<size_t>(P * N), 0.0f);
+    for (int64_t m = 0; m < M; ++m)
+      for (int64_t n = 0; n < N; ++n) {
+        float acc = 0.0f;
+        for (int64_t k = 0; k < K; ++k)
+          acc += act_r[static_cast<size_t>(m * K + k)] * w_f32[static_cast<size_t>(n * K + k)];
+        ref[static_cast<size_t>(m * N + n)] = acc;
+      }
+
+    // Repack weight + process E8M0 scales (mxfp4: passthrough permute, no global).
+    DeviceTensor dp(gpu, gq.q, DType::kI8, {N, K / 2}, w.packed.data());
+    DeviceTensor ds(gpu, gq.q, DType::kI8, {N, K / 32}, w.scale.data());
+    DeviceTensor wq(gpu, gq.q, DType::kI32, {E, K / 16, N * 2});
+    DeviceTensor sc(gpu, gq.q, DType::kI8, {E, K / 32, N});  // K/32 groups (group_blocks=2)
+    vt::cuda::MarlinRepackExpertWeight(stream, dev, static_cast<uint32_t*>(wq.ptr()),
+                                       static_cast<const uint8_t*>(dp.ptr()),
+                                       static_cast<int>(K), static_cast<int>(N));
+    vt::cuda::MarlinProcessExpertScalesMxfp4(stream, static_cast<const uint8_t*>(ds.ptr()),
+                                             static_cast<uint8_t*>(sc.ptr()),
+                                             static_cast<int>(K), static_cast<int>(N));
+    float g_dummy = 1.0f;  // ignored on the mxfp4 path (kernel skips global for E8M0)
+    DeviceTensor gg(gpu, gq.q, DType::kF32, {E}, &g_dummy);
+
+    // Align inputs: all P tokens -> expert 0.
+    std::vector<int32_t> topk_ids(static_cast<size_t>(P), 0);
+    std::vector<float> topk_w(static_cast<size_t>(P), 1.0f);
+    const int block = vt::cuda::MarlinMoeAlignBlockSizeSelect(static_cast<int>(M),
+                                                              static_cast<int>(top_k),
+                                                              static_cast<int>(E));
+    int max_tok = 0, max_blk = 0;
+    vt::cuda::MarlinMoeAlignSizes(static_cast<int>(M), static_cast<int>(top_k),
+                                  static_cast<int>(E), block, &max_tok, &max_blk);
+    DeviceTensor dtid(gpu, gq.q, DType::kI32, {M, top_k}, topk_ids.data());
+    DeviceTensor dtw(gpu, gq.q, DType::kF32, {M, top_k}, topk_w.data());
+    DeviceTensor sorted_ids(gpu, gq.q, DType::kI32, {max_tok});
+    DeviceTensor expert_ids(gpu, gq.q, DType::kI32, {max_blk});
+    DeviceTensor num_pad(gpu, gq.q, DType::kI32, {1});
+    vt::cuda::MarlinMoeAlignBlockSize(stream, static_cast<const int32_t*>(dtid.ptr()),
+                                      static_cast<int>(M), static_cast<int>(top_k),
+                                      static_cast<int>(E), block,
+                                      static_cast<int32_t*>(sorted_ids.ptr()),
+                                      static_cast<int32_t*>(expert_ids.ptr()),
+                                      static_cast<int32_t*>(num_pad.ptr()));
+    const int sms = vt::cuda::MarlinDeviceSms(dev);
+    DeviceTensor ws(gpu, gq.q, DType::kI32, {sms * 4});
+    gpu.Memset(gq.q, ws.ptr(), 0, static_cast<size_t>(sms) * 4 * sizeof(int32_t));
+    DeviceTensor dact(gpu, gq.q, DType::kBF16, {M, K}, act_bf16.data());
+    DeviceTensor dout(gpu, gq.q, DType::kBF16, {P, N});
+
+    vt::MoeMarlinArgs args{block, static_cast<int>(top_k), static_cast<int>(M),
+                           static_cast<int>(N), static_cast<int>(K), false};
+    args.group_size = 32;
+    args.mxfp4 = true;
+    vt::MoeGroupedGemmNvfp4Marlin(gq.q, dout.tensor(), dact.tensor(), wq.tensor(), sc.tensor(),
+                                  gg.tensor(), ws.tensor(), sorted_ids.tensor(),
+                                  expert_ids.tensor(), num_pad.tensor(), dtw.tensor(), args);
+    std::vector<uint16_t> got_bf16(static_cast<size_t>(P * N));
+    dout.Download(gq.q, got_bf16.data());
+    std::vector<float> got(got_bf16.size());
+    for (size_t i = 0; i < got.size(); ++i) got[i] = vt::BF16ToF32(got_bf16[i]);
+    // Report the true error magnitude: a correct bf16-out Marlin GEMM vs an f32
+    // reference sits near bf16 rounding (~1e-2); a systematic error compounds.
+    float max_rel = 0.0f, max_abs = 0.0f;
+    for (size_t i = 0; i < got.size(); ++i) {
+      const float a = std::fabs(got[i] - ref[i]);
+      const float r = a / (std::fabs(ref[i]) + 1e-6f);
+      if (a > max_abs) max_abs = a;
+      if (r > max_rel) max_rel = r;
+    }
+    MESSAGE("MXFP4 K=" << K << " N=" << N << " M=" << M
+                       << " max_abs=" << max_abs << " max_rel=" << max_rel);
+    // Measured max_rel ~= 3.8e-3 at M=1,M=8 (K=256,N=128; pure bf16 rounding, NOT
+    // a systematic error). Gate at 2e-2 leaves bf16 headroom. The added Qwen3-8B
+    // shapes ({4096,4096},{4096,12288},{12288,4096}) are RUN-PENDING (box contended
+    // when authored) — a FAIL there localizes a large-N/K group_blocks=2 kernel bug
+    // as the e2e residual's cause; a PASS shifts the residual to model integration.
+    CheckClose(got, ref, 2e-2f, 2e-2f);
+  }
+  }
+}
+
 TEST_CASE("CUDA moe_align parallel == serial (expert_ids/num_pad exact, per-expert multiset)") {
   if (!HasCuda()) {
     MESSAGE("no CUDA backend registered; skipping");