diff --git a/docs/kernel-contract.md b/docs/kernel-contract.md
index e625bcb2fd..0b80446a42 100644
--- a/docs/kernel-contract.md
+++ b/docs/kernel-contract.md
@@ -108,3 +108,7 @@ CUDA: the RMS norm rows (16-byte form and the fp32-out tree), Nemotron's router
 tile, lane_gemv and its split-K form (both bit-equal to it), prompt16's FP4 tile with bf16 and fp32 out, qmmf's FP4
 tiles with fp32 out (K slices met in a cluster, fused, or by `reduce_kernel<true>`), and the reduce-scatter's
 `tf_rank_sum`.
+
+## CUDA compiler resource baseline
+
+The registered usage arrays were re-derived with NVCC 13.3.73 (`cuda_13.3.r13.3/compiler.38244171_0`) for SM 80, 86, 89, 120 and 121. `kernel-contract-limits` reads each compiled image and reports its register count, static shared memory and thread bound. Source changes or a different compiler may change these values; re-derive the arrays and run the live audit rather than changing the audit tolerance. Launch dimensions, dynamic shared memory, device limits and planted controls remain independent checks. Only SM 120 was executed on hardware in this update; the other four records come from their compiled images, not hardware execution.
diff --git a/zig/build/cuda.zig b/zig/build/cuda.zig
index 684d6c3cc9..8aee541747 100644
--- a/zig/build/cuda.zig
+++ b/zig/build/cuda.zig
@@ -240,6 +240,18 @@ pub fn hostTests(b: *std.Build, draft_ids: *std.Build.Module, build_options: *st
     step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = b.createModule(.{ .root_source_file = b.path("zig/build/check_fatbins.zig"), .target = host }) })).step);
     const cuda = runtime(b, host, .debug, &.{});
     const mods = family(b, host, .debug, cuda, draft_ids);
+    const reference_step = b.step("test-cuda-reference", "CUDA's FP64 comparison with CPU-only readback (no GPU)");
+    for ([_]std.builtin.OptimizeMode{ .debug, .fast }) |mode| {
+        const reference = b.addRunArtifact(b.addTest(.{ .root_module = b.createModule(.{
+            .root_source_file = b.path("zig/tests/cuda/quant_ref_host_tests.zig"),
+            .target = host,
+            .optimize = mode,
+            .link_libc = true,
+            .imports = &.{ .{ .name = "cuda", .module = cuda }, .{ .name = "core", .module = mods.core } },
+        }) }));
+        step.dependOn(&reference.step);
+        reference_step.dependOn(&reference.step);
+    }
     const native_modules = engines(b, host, .debug, cuda, mods.core, mods.lanes, mods.nemotron, mods.heat);
     const native = native_modules.engines;
     // the HTTP server's unit tests (zig/src/server/root.zig), which the macOS build runs, on Linux hosts too
diff --git a/zig/kernels/cuda/qmmf.cu b/zig/kernels/cuda/qmmf.cu
index 65cf08fd22..8af69d3b33 100644
--- a/zig/kernels/cuda/qmmf.cu
+++ b/zig/kernels/cuda/qmmf.cu
@@ -32,6 +32,11 @@ __device__ __forceinline__ uint32_t fp8pair(uint32_t w) {
     return r;
 }
 
+// UE8M0 has no zero encoding: byte 0 is 2^-127, byte 255 is NaN.
+__device__ __forceinline__ float e8m0f(uint8_t b) {
+    return __uint_as_float(b == 0 ? 0x00400000u : b == 255 ? 0x7FC00000u : uint32_t(b) << 23);
+}
+
 __device__ __forceinline__ uint32_t comp(const uint4& v, int c) { return c == 0 ? v.x : c == 1 ? v.y : c == 2 ? v.z : v.w; }
 
 __device__ __forceinline__ float e4m3f(uint8_t b) {
@@ -193,8 +198,8 @@ __global__ void __launch_bounds__(WM * WN * 32) qmmf_kernel(
                             s0 = e4m3f(static_cast<uint8_t>(sq[j][0] >> (8 * blk)));
                             s1 = e4m3f(static_cast<uint8_t>(sq[j][1] >> (8 * blk)));
                         } else {
-                            s0 = __int_as_float(static_cast<int>((sq[j][0] >> (8 * blk)) & 0xFFu) << 23);
-                            s1 = __int_as_float(static_cast<int>((sq[j][1] >> (8 * blk)) & 0xFFu) << 23);
+                            s0 = e8m0f(static_cast<uint8_t>(sq[j][0] >> (8 * blk)));
+                            s1 = e8m0f(static_cast<uint8_t>(sq[j][1] >> (8 * blk)));
                         }
 #pragma unroll
                         for (int i = 0; i < T::MT; ++i)
@@ -293,6 +298,14 @@ __global__ void reduce_kernel(const float* __restrict__ part, void* __restrict__
     else reinterpret_cast<__nv_bfloat16*>(out)[i] = __float2bfloat16_rn(acc);
 }
 
+template __global__ void qmmf_kernel<MXFP8, 16, 64, 1, 4, 4, false, false, false>(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int);
+template __global__ void qmmf_kernel<MXFP8, 16, 64, 1, 4, 4, false, true, false>(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int);
+template __global__ void qmmf_kernel<MXFP8, 32, 64, 1, 4, 4, false, false, false>(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int);
+template __global__ void qmmf_kernel<MXFP8, 32, 64, 1, 4, 4, false, true, false>(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int);
+template __global__ void qmmf_kernel<MXFP8, 64, 64, 1, 4, 4, false, false, false>(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int);
+template __global__ void qmmf_kernel<MXFP8, 64, 64, 1, 4, 4, false, true, false>(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int);
+template __global__ void qmmf_kernel<MXFP8, 64, 64, 1, 4, 4, false, false, true>(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int);
+
 // The instances the Zig launchers load by name (the Python host instantiates these through its launch switch).
 template __global__ void qmmf_kernel<FP8G, 16, 64, 1, 4, 4, false, false, false>(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int);
 template __global__ void qmmf_kernel<FP8G, 16, 64, 1, 4, 4, false, true, false>(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int);
diff --git a/zig/src/core/contract/registry.zig b/zig/src/core/contract/registry.zig
index b79337b5c5..555dc07840 100644
--- a/zig/src/core/contract/registry.zig
+++ b/zig/src/core/contract/registry.zig
@@ -20,23 +20,31 @@ const full72_symbol = "_ZN20tf_prefill_attention11full_kernelILi72ELi8ELi8EEEvPK
 /// CUDA kernels as launched; tf-cuda-test kernel-contract-limits re-derives the usage records.
 pub const cuda = [_]CudaRecord{
     .{ .launch = .{ .name = "norms tf_rms_rows_bf16", .symbol = "tf_rms_rows_bf16", .image = "norms", .threads = 256 }, .usage = .{ .{ .regs = 79, .static = 32, .bound = 256 }, .{ .regs = 79, .static = 32, .bound = 256 }, .{ .regs = 79, .static = 32, .bound = 256 }, .{ .regs = 80, .static = 32, .bound = 256 }, .{ .regs = 80, .static = 32, .bound = 256 } } },
-    .{ .launch = .{ .name = "norms tf_rms_rows", .symbol = "tf_rms_rows", .image = "norms", .threads = 1024 }, .usage = .{ .{ .regs = 32, .static = 132, .bound = 1024 }, .{ .regs = 46, .static = 132, .bound = 1024 }, .{ .regs = 46, .static = 132, .bound = 1024 }, .{ .regs = 45, .static = 132, .bound = 1024 }, .{ .regs = 45, .static = 132, .bound = 1024 } } },
+    .{ .launch = .{ .name = "norms tf_rms_rows", .symbol = "tf_rms_rows", .image = "norms", .threads = 1024 }, .usage = .{ .{ .regs = 32, .static = 132, .bound = 1024 }, .{ .regs = 46, .static = 132, .bound = 1024 }, .{ .regs = 46, .static = 132, .bound = 1024 }, .{ .regs = 43, .static = 132, .bound = 1024 }, .{ .regs = 43, .static = 132, .bound = 1024 } } },
     .{ .launch = .{ .name = "router tf_nemo_router", .symbol = "tf_nemo_router", .image = "nemotron_route", .threads = 256, .dynamic = 28800 }, .usage = .{ .{ .regs = 32, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 } } },
-    .{ .launch = .{ .name = "router tf_nemo_topk", .symbol = "tf_nemo_topk", .image = "nemotron_route", .threads = 32 }, .usage = .{ .{ .regs = 94, .static = 0, .bound = 32 }, .{ .regs = 95, .static = 0, .bound = 32 }, .{ .regs = 95, .static = 0, .bound = 32 }, .{ .regs = 91, .static = 0, .bound = 32 }, .{ .regs = 91, .static = 0, .bound = 32 } } },
+    .{ .launch = .{ .name = "router tf_nemo_topk", .symbol = "tf_nemo_topk", .image = "nemotron_route", .threads = 32 }, .usage = .{ .{ .regs = 94, .static = 0, .bound = 32 }, .{ .regs = 95, .static = 0, .bound = 32 }, .{ .regs = 95, .static = 0, .bound = 32 }, .{ .regs = 96, .static = 0, .bound = 32 }, .{ .regs = 96, .static = 0, .bound = 32 } } },
     .{ .launch = .{ .name = "qmm_group cluster of 8", .symbol = group_symbol, .image = "qmm_group", .threads = 128, .dynamic = 35328, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 54, .static = 0, .bound = 128 }, .{ .regs = 54, .static = 0, .bound = 128 }, .{ .regs = 54, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 } } },
     .{ .launch = .{ .name = "qmm_group one CTA a tile", .symbol = group_symbol, .image = "qmm_group", .threads = 128, .dynamic = 35328 }, .usage = .{ .{ .regs = 54, .static = 0, .bound = 128 }, .{ .regs = 54, .static = 0, .bound = 128 }, .{ .regs = 54, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 } } },
-    .{ .launch = .{ .name = "lane_gemv", .symbol = gemv_symbol, .image = "lane_gemv", .threads = 128, .dynamic = 35328 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 57, .static = 0, .bound = 128 }, .{ .regs = 57, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 } } },
-    .{ .launch = .{ .name = "lane_gemv split-K", .symbol = "tf_lane_gemv_split", .image = "lane_gemv", .threads = 128, .dynamic = 17664 }, .usage = .{ .{ .regs = 62, .static = 128, .bound = 128 }, .{ .regs = 56, .static = 128, .bound = 128 }, .{ .regs = 56, .static = 128, .bound = 128 }, .{ .regs = 44, .static = 128, .bound = 128 }, .{ .regs = 44, .static = 128, .bound = 128 } } },
+    .{ .launch = .{ .name = "lane_gemv", .symbol = gemv_symbol, .image = "lane_gemv", .threads = 128, .dynamic = 35328 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 58, .static = 0, .bound = 128 }, .{ .regs = 58, .static = 0, .bound = 128 }, .{ .regs = 66, .static = 0, .bound = 128 }, .{ .regs = 66, .static = 0, .bound = 128 } } },
+    .{ .launch = .{ .name = "lane_gemv split-K", .symbol = "tf_lane_gemv_split", .image = "lane_gemv", .threads = 128, .dynamic = 17664 }, .usage = .{ .{ .regs = 62, .static = 128, .bound = 128 }, .{ .regs = 56, .static = 128, .bound = 128 }, .{ .regs = 56, .static = 128, .bound = 128 }, .{ .regs = 72, .static = 128, .bound = 128 }, .{ .regs = 72, .static = 128, .bound = 128 } } },
     .{ .launch = .{ .name = "prompt16 FP4", .symbol = std.fmt.comptimePrint(prompt_symbol, .{ 0, 0 }), .image = "prompt16", .threads = 128, .dynamic = 41984 }, .usage = .{ .{ .regs = 255, .static = 0, .bound = 128 }, .{ .regs = 255, .static = 0, .bound = 128 }, .{ .regs = 255, .static = 0, .bound = 128 }, .{ .regs = 212, .static = 0, .bound = 128 }, .{ .regs = 212, .static = 0, .bound = 128 } } },
-    .{ .launch = .{ .name = "prompt16 FP8", .symbol = std.fmt.comptimePrint(prompt_symbol, .{ 1, 0 }), .image = "prompt16", .threads = 128, .dynamic = 49152 }, .usage = .{ .{ .regs = 243, .static = 0, .bound = 128 }, .{ .regs = 243, .static = 0, .bound = 128 }, .{ .regs = 243, .static = 0, .bound = 128 }, .{ .regs = 218, .static = 0, .bound = 128 }, .{ .regs = 218, .static = 0, .bound = 128 } } },
+    .{ .launch = .{ .name = "prompt16 FP8", .symbol = std.fmt.comptimePrint(prompt_symbol, .{ 1, 0 }), .image = "prompt16", .threads = 128, .dynamic = 49152 }, .usage = .{ .{ .regs = 243, .static = 0, .bound = 128 }, .{ .regs = 243, .static = 0, .bound = 128 }, .{ .regs = 243, .static = 0, .bound = 128 }, .{ .regs = 212, .static = 0, .bound = 128 }, .{ .regs = 212, .static = 0, .bound = 128 } } },
     .{ .launch = .{ .name = "prompt16 FP4 fp32 out", .symbol = std.fmt.comptimePrint(prompt_symbol, .{ 0, 1 }), .image = "prompt16", .threads = 128, .dynamic = 41984 }, .usage = .{ .{ .regs = 255, .static = 0, .bound = 128 }, .{ .regs = 255, .static = 0, .bound = 128 }, .{ .regs = 255, .static = 0, .bound = 128 }, .{ .regs = 212, .static = 0, .bound = 128 }, .{ .regs = 212, .static = 0, .bound = 128 } } },
     .{ .launch = .{ .name = "qmmf FP4 fp32 out 16-row tile", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 16, 0, 0 }), .image = "qmmf", .threads = 128, .dynamic = 17408 }, .usage = .{ .{ .regs = 88, .static = 0, .bound = 128 }, .{ .regs = 88, .static = 0, .bound = 128 }, .{ .regs = 96, .static = 0, .bound = 128 }, .{ .regs = 56, .static = 0, .bound = 128 }, .{ .regs = 56, .static = 0, .bound = 128 } } },
     .{ .launch = .{ .name = "qmmf FP4 fp32 out 16-row tile, cluster of 8", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 16, 1, 0 }), .image = "qmmf", .threads = 128, .dynamic = 17408, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 } } },
     .{ .launch = .{ .name = "qmmf FP4 fp32 out 32-row tile", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 32, 0, 0 }), .image = "qmmf", .threads = 128, .dynamic = 25600 }, .usage = .{ .{ .regs = 106, .static = 0, .bound = 128 }, .{ .regs = 106, .static = 0, .bound = 128 }, .{ .regs = 128, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 } } },
-    .{ .launch = .{ .name = "qmmf FP4 fp32 out 32-row tile, cluster of 8", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 32, 1, 0 }), .image = "qmmf", .threads = 128, .dynamic = 25600, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 80, .static = 0, .bound = 128 }, .{ .regs = 63, .static = 0, .bound = 128 }, .{ .regs = 63, .static = 0, .bound = 128 } } },
+    .{ .launch = .{ .name = "qmmf FP4 fp32 out 32-row tile, cluster of 8", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 32, 1, 0 }), .image = "qmmf", .threads = 128, .dynamic = 25600, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 80, .static = 0, .bound = 128 }, .{ .regs = 61, .static = 0, .bound = 128 }, .{ .regs = 61, .static = 0, .bound = 128 } } },
     .{ .launch = .{ .name = "qmmf FP4 fp32 out 64-row tile", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 64, 0, 0 }), .image = "qmmf", .threads = 128, .dynamic = 41984 }, .usage = .{ .{ .regs = 148, .static = 0, .bound = 128 }, .{ .regs = 148, .static = 0, .bound = 128 }, .{ .regs = 180, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 } } },
     .{ .launch = .{ .name = "qmmf FP4 fp32 out 64-row tile, cluster of 8", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 64, 1, 0 }), .image = "qmmf", .threads = 128, .dynamic = 41984, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 72, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 } } },
     .{ .launch = .{ .name = "qmmf FP4 fp32 out fused slices", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 64, 0, 1 }), .image = "qmmf", .threads = 128, .dynamic = 41984 }, .usage = .{ .{ .regs = 182, .static = 0, .bound = 128 }, .{ .regs = 182, .static = 0, .bound = 128 }, .{ .regs = 228, .static = 0, .bound = 128 }, .{ .regs = 168, .static = 0, .bound = 128 }, .{ .regs = 168, .static = 0, .bound = 128 } } },
+    .{ .launch = .{ .name = "qmmf MXFP8 16-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi16ELi64ELi1ELi4ELi4ELb0ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 25088 }, .usage = .{ .{ .regs = 94, .static = 0, .bound = 128 }, .{ .regs = 94, .static = 0, .bound = 128 }, .{ .regs = 94, .static = 0, .bound = 128 }, .{ .regs = 60, .static = 0, .bound = 128 }, .{ .regs = 60, .static = 0, .bound = 128 } } },
+    .{ .launch = .{ .name = "qmmf MXFP8 16-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi16ELi64ELi1ELi4ELi4ELb0ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 25088, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 69, .static = 0, .bound = 128 }, .{ .regs = 69, .static = 0, .bound = 128 }, .{ .regs = 69, .static = 0, .bound = 128 }, .{ .regs = 56, .static = 0, .bound = 128 }, .{ .regs = 56, .static = 0, .bound = 128 } } },
+    .{ .launch = .{ .name = "qmmf MXFP8 32-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi32ELi64ELi1ELi4ELi4ELb0ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 33280 }, .usage = .{ .{ .regs = 128, .static = 0, .bound = 128 }, .{ .regs = 128, .static = 0, .bound = 128 }, .{ .regs = 128, .static = 0, .bound = 128 }, .{ .regs = 57, .static = 0, .bound = 128 }, .{ .regs = 57, .static = 0, .bound = 128 } } },
+    .{ .launch = .{ .name = "qmmf MXFP8 32-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi32ELi64ELi1ELi4ELi4ELb0ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 33280, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 80, .static = 0, .bound = 128 }, .{ .regs = 80, .static = 0, .bound = 128 }, .{ .regs = 80, .static = 0, .bound = 128 }, .{ .regs = 66, .static = 0, .bound = 128 }, .{ .regs = 66, .static = 0, .bound = 128 } } },
+    .{ .launch = .{ .name = "qmmf MXFP8 64-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi64ELi64ELi1ELi4ELi4ELb0ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 49664, .opt_in = true }, .usage = .{ .{ .regs = 168, .static = 0, .bound = 128 }, .{ .regs = 168, .static = 0, .bound = 128 }, .{ .regs = 168, .static = 0, .bound = 128 }, .{ .regs = 108, .static = 0, .bound = 128 }, .{ .regs = 108, .static = 0, .bound = 128 } } },
+    .{ .launch = .{ .name = "qmmf MXFP8 64-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi64ELi64ELi1ELi4ELi4ELb0ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 49664, .opt_in = true, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 72, .static = 0, .bound = 128 }, .{ .regs = 72, .static = 0, .bound = 128 }, .{ .regs = 72, .static = 0, .bound = 128 }, .{ .regs = 103, .static = 0, .bound = 128 }, .{ .regs = 103, .static = 0, .bound = 128 } } },
+    .{ .launch = .{ .name = "qmmf MXFP8 64-row tile, fused slices", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi64ELi64ELi1ELi4ELi4ELb0ELb0ELb1EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 49664, .opt_in = true }, .usage = .{ .{ .regs = 229, .static = 0, .bound = 128 }, .{ .regs = 229, .static = 0, .bound = 128 }, .{ .regs = 229, .static = 0, .bound = 128 }, .{ .regs = 158, .static = 0, .bound = 128 }, .{ .regs = 158, .static = 0, .bound = 128 } } },
+    .{ .launch = .{ .name = "qmmf reduce_kernel<false>", .symbol = "_ZN7tf_qmmf13reduce_kernelILb0EEEvPKfPvxif", .image = "qmmf", .threads = 256 }, .usage = .{ .{ .regs = 32, .static = 0, .bound = 0 }, .{ .regs = 36, .static = 0, .bound = 0 }, .{ .regs = 36, .static = 0, .bound = 0 }, .{ .regs = 40, .static = 0, .bound = 0 }, .{ .regs = 40, .static = 0, .bound = 0 } } },
     .{ .launch = .{ .name = "qmmf reduce_kernel<true>", .symbol = "_ZN7tf_qmmf13reduce_kernelILb1EEEvPKfPvxif", .image = "qmmf", .threads = 256 }, .usage = .{ .{ .regs = 32, .static = 0, .bound = 0 }, .{ .regs = 36, .static = 0, .bound = 0 }, .{ .regs = 36, .static = 0, .bound = 0 }, .{ .regs = 40, .static = 0, .bound = 0 }, .{ .regs = 40, .static = 0, .bound = 0 } } },
     .{ .launch = .{ .name = "rank_sum tf_rank_sum", .symbol = "tf_rank_sum", .image = "rank_sum", .threads = 256 }, .usage = .{ .{ .regs = 32, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 } } },
     .{ .launch = .{ .name = "vit vision_kernel<80, 8> (RADIO)", .symbol = vision80_symbol, .image = "prefill_attention", .threads = 256, .dynamic = 45056 }, .usage = .{ .{ .regs = 143, .static = 0, .bound = 256 }, .{ .regs = 143, .static = 0, .bound = 256 }, .{ .regs = 143, .static = 0, .bound = 256 }, .{ .regs = 122, .static = 0, .bound = 256 }, .{ .regs = 122, .static = 0, .bound = 256 } } },
diff --git a/zig/src/cuda/mxfp8.md b/zig/src/cuda/mxfp8.md
new file mode 100644
index 0000000000..5878df43eb
--- /dev/null
+++ b/zig/src/cuda/mxfp8.md
@@ -0,0 +1,27 @@
+# Native MXFP8 projections
+
+`cuda.mxfp8` packs row-major E4M3 code bytes and one UE8M0 scale byte per
+32 input channels for `qlinear.Weight.mxfp8`. Both prompt and decode calls use
+the lane matmul, with BF16 inputs, FP32 accumulation and BF16 output. Packing
+permutes bytes; it does not requantize weights or expand them to BF16.
+
+The logical shape is `[N, K]`, with positive dimensions and K divisible by 64.
+`packCodes` writes `padded(N) * K` bytes; `packScales` writes
+`padded(N) * K/32` bytes. Rows are padded to 128; padded codes are zero and
+padded scales are one. `weight` returns the borrowed device view. The caller
+owns its device buffers and the lane's split-K scratch.
+
+UE8M0 bytes 0 through 254 mean `2^(byte - 127)`; byte 255 is NaN. In
+particular, byte zero is a nonzero subnormal FP32 scale. See the
+[CUDA datatype contract](https://docs.nvidia.com/cuda/archive/12.8.0/cuda-math-api/cuda_math_api/struct____nv__fp8__e8m0.html).
+
+`tf-cuda-test mxfp8` checks all 256 scale encodings through the production
+matmul, then compares synthetic projections with an independent FP64
+calculation. It also checks strided activations, partial output tiles,
+row counts 1/7/17/33/256, prompt/decode dispatch, and cluster versus explicit
+split reduction. Host tests cover scale permutation, padding, shared-memory
+sizes, symbols and format refusals.
+
+This primitive does not admit new checkpoint families or certify model
+correctness, tensor parallelism or performance. Rebuild the qmmf fatbin with
+the native runtime; previous fatbins lack the MXFP8 instantiations.
diff --git a/zig/src/cuda/mxfp8.zig b/zig/src/cuda/mxfp8.zig
new file mode 100644
index 0000000000..7c57aa50df
--- /dev/null
+++ b/zig/src/cuda/mxfp8.zig
@@ -0,0 +1,41 @@
+//! E4M3 codes with one UE8M0 byte per 32 inputs, repacked losslessly for the lane matmul.
+const std = @import("std");
+const fp8 = @import("fp8.zig");
+const qmmf = @import("qmmf.zig");
+
+pub const group = 32;
+pub const padded = fp8.padded;
+pub const packCodes = fp8.packCodes;
+
+/// Row-major [n, k/32] bytes to [npad/64][k/64][64][2]; padded rows have scale 1.
+pub fn packScales(out: []u8, scales: []const u8, n: usize, k: usize) void {
+    const npad = padded(n);
+    const kg = k / qmmf.group;
+    std.debug.assert(k % qmmf.group == 0 and out.len == npad * (k / group) and scales.len == n * (k / group));
+    for (0..npad / 64) |t| for (0..kg) |g| for (0..64) |r| for (0..2) |h| {
+        const row = t * 64 + r;
+        out[((t * kg + g) * 64 + r) * 2 + h] = if (row < n) scales[row * (k / group) + g * 2 + h] else 127;
+    };
+}
+
+pub fn weight(codes: u64, scales: u64, n: usize, k: usize) qmmf.Weight {
+    return .{ .mode = .mxfp8, .codes = codes, .scales = scales, .n = @intCast(n), .k = @intCast(k), .npad = @intCast(padded(n)) };
+}
+
+test "every scale byte lands in its independent 32-input group, including padded rows" {
+    const a = std.testing.allocator;
+    const n = 130;
+    const k = 256;
+    const scales = try a.alloc(u8, n * (k / group));
+    defer a.free(scales);
+    for (scales, 0..) |*v, i| v.* = @truncate(i * 13 + 7);
+    const out = try a.alloc(u8, padded(n) * (k / group));
+    defer a.free(out);
+    packScales(out, scales, n, k);
+    for (0..padded(n)) |row| for (0..k / group) |g| {
+        const at = (((row / 64) * (k / 64) + g / 2) * 64 + row % 64) * 2 + g % 2;
+        try std.testing.expectEqual(if (row < n) scales[row * (k / group) + g] else @as(u8, 127), out[at]);
+    };
+    const w = weight(1, 2, n, k);
+    try std.testing.expect(w.mode == .mxfp8 and w.npad == 256 and w.scale == 1.0);
+}
diff --git a/zig/src/cuda/qlinear.zig b/zig/src/cuda/qlinear.zig
index dbb86bb8d0..a5031b6bcd 100644
--- a/zig/src/cuda/qlinear.zig
+++ b/zig/src/cuda/qlinear.zig
@@ -34,6 +34,7 @@ pub const Affine4 = struct {
 pub const Weight = union(enum) {
     affine4: Affine4,
     fp8g: qmmf.Weight,
+    mxfp8: qmmf.Weight, // e4m3 with UE8M0 scale bytes per 32 inputs
     nvfp4: qmmf.Weight,
     fp8: qmmf.Weight, // e4m3 with one fp32 scale (ModelOpt FP8)
     bf16: qmmf.Weight, // unquantized: the checkpoint's bf16 values
@@ -41,14 +42,14 @@ pub const Weight = union(enum) {
     pub fn outputs(w: Weight) usize {
         return switch (w) {
             .affine4 => |q| q.n,
-            .fp8g, .nvfp4, .fp8, .bf16 => |q| q.n,
+            .fp8g, .mxfp8, .nvfp4, .fp8, .bf16 => |q| q.n,
         };
     }
 
     pub fn inputs(w: Weight) usize {
         return switch (w) {
             .affine4 => |q| q.k,
-            .fp8g, .nvfp4, .fp8, .bf16 => |q| q.k,
+            .fp8g, .mxfp8, .nvfp4, .fp8, .bf16 => |q| q.k,
         };
     }
 };
@@ -81,6 +82,7 @@ pub const Linear = struct {
                 return a.decode(s, in.x, in.sums, q, out, rows);
             },
             .fp8g => |q| return l.onLane(s, q, .fp8g, in, out, rows),
+            .mxfp8 => |q| return l.onLane(s, q, .mxfp8, in, out, rows),
             .nvfp4 => |q| return l.onLane(s, q, .fp4, in, out, rows),
             .fp8 => |q| return if (q.act != 0 and l.ck != null) l.checkpoint(s, l.ck.?, q, in, out, rows) else l.onLane(s, q, .fp8, in, out, rows),
             .bf16 => |q| return l.onLane(s, q, .bf16, in, out, rows),
@@ -96,6 +98,7 @@ pub const Linear = struct {
                 return a.prompt(s, in.x, q, out, rows);
             },
             .fp8g => |q| return l.onLane(s, q, .fp8g, in, out, rows),
+            .mxfp8 => |q| return l.onLane(s, q, .mxfp8, in, out, rows),
             .nvfp4 => |q| return l.onPrompt(s, q, .fp4, in, out, rows),
             .fp8 => |q| return if (q.act != 0 and l.ck != null) l.checkpoint(s, l.ck.?, q, in, out, rows) else l.onPrompt(s, q, .fp8, in, out, rows),
             .bf16 => |q| return l.onLane(s, q, .bf16, in, out, rows),
@@ -385,6 +388,8 @@ test "a lane weight whose mode is not its format's is refused before any launch"
     const fp4: qmmf.Weight = .{ .mode = .fp4, .codes = 0, .scales = 0, .n = 64, .k = 64, .npad = 128 };
     const s: Stream = undefined;
     try std.testing.expectError(error.FormatMismatch, l.decode(s, .{ .fp8g = fp4 }, .{ .x = 0 }, 0, 1));
+    try std.testing.expectError(error.FormatMismatch, l.decode(s, .{ .mxfp8 = fp4 }, .{ .x = 0 }, 0, 1));
+    try std.testing.expectError(error.FormatMismatch, l.prompt(s, .{ .mxfp8 = fp4 }, .{ .x = 0 }, 0, 1));
     try std.testing.expectError(error.FormatNotLoaded, l.prompt(s, .{ .nvfp4 = fp4 }, .{ .x = 0 }, 0, 1));
     const f8: qmmf.Weight = .{ .mode = .fp8g, .codes = 0, .scales = 0, .n = 64, .k = 64, .npad = 128 };
     try std.testing.expectError(error.EpilogueNotBuilt, l.decodeF32(s, .{ .fp8g = f8 }, .{ .x = 0 }, 0, 1));
diff --git a/zig/src/cuda/qmmf.zig b/zig/src/cuda/qmmf.zig
index ec7fcd6e04..da5f1e0e92 100644
--- a/zig/src/cuda/qmmf.zig
+++ b/zig/src/cuda/qmmf.zig
@@ -1,4 +1,4 @@
-//! The lane matmul (qmmf.cu) for FP4, FP8, FP8G and BF16 weights, with its kernel, tile and slices.
+//! The lane matmul (qmmf.cu) for FP4, FP8, MXFP8, FP8G and BF16 weights, with its kernel, tile and slices.
 const std = @import("std");
 const driver = @import("driver.zig");
 const module = @import("module.zig");
@@ -17,9 +17,9 @@ const bn = 64;
 const threads = 128; // WM 1 x WN 4 warps
 
 /// The packed formats (qmmf.cu Mode): FP4 with e4m3 scales a 16, e4m3 with one scale or fp32 a 64, and bf16.
-pub const Mode = enum(u3) { fp4 = 0, fp8 = 1, fp8g = 3, bf16 = 4 };
+pub const Mode = enum(u3) { fp4 = 0, fp8 = 1, mxfp8 = 2, fp8g = 3, bf16 = 4 };
 
-const modes = [_]Mode{ .fp4, .fp8, .fp8g, .bf16 };
+const modes = [_]Mode{ .fp4, .fp8, .mxfp8, .fp8g, .bf16 };
 
 fn modeIndex(mode: Mode) usize {
     return std.mem.indexOfScalar(Mode, &modes, mode).?;
@@ -29,10 +29,10 @@ fn modeIndex(mode: Mode) usize {
 pub fn sharedBytes(mode: Mode, bm: u32) u32 {
     const w: u32 = switch (mode) {
         .fp4 => bn * group / 2,
-        .fp8, .fp8g => bn * group,
+        .fp8, .mxfp8, .fp8g => bn * group,
         .bf16 => bn * group * 2,
     };
-    const s: u32 = if (mode == .fp4 or mode == .fp8g) bn * 4 else 0;
+    const s: u32 = if (mode == .fp4 or mode == .fp8g) bn * 4 else if (mode == .mxfp8) bn * 2 else 0;
     const stage = (bm * group * 2 + w + s + 127) / 128 * 128;
     const partials = (bm / 16) * 2 * 4 * threads * 4;
     return @max(4 * stage, partials);
@@ -220,6 +220,9 @@ test "tile shared bytes, slices and row tiles as the Python host picks them" {
     try std.testing.expectEqual(@as(u32, 25600), sharedBytes(.fp8g, 16));
     try std.testing.expectEqual(@as(u32, 33792), sharedBytes(.fp8g, 32));
     try std.testing.expectEqual(@as(u32, 50176), sharedBytes(.fp8g, 64));
+    try std.testing.expectEqual(@as(u32, 25088), sharedBytes(.mxfp8, 16));
+    try std.testing.expectEqual(@as(u32, 33280), sharedBytes(.mxfp8, 32));
+    try std.testing.expectEqual(@as(u32, 49664), sharedBytes(.mxfp8, 64));
     try std.testing.expectEqual(@as(u32, 17408), sharedBytes(.fp4, 16));
     try std.testing.expectEqual(@as(u32, 25600), sharedBytes(.fp4, 32));
     try std.testing.expectEqual(@as(u32, 41984), sharedBytes(.fp4, 64));
@@ -243,6 +246,10 @@ test "tile shared bytes, slices and row tiles as the Python host picks them" {
 }
 
 test "instance names are qmmf.cu's" {
+    try std.testing.expectEqualStrings(
+        "_ZN7tf_qmmf11qmmf_kernelILi2ELi16ELi64ELi1ELi4ELi4ELb0ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii",
+        symbol(.mxfp8, 16, true, false, false),
+    );
     try std.testing.expectEqualStrings(
         "_ZN7tf_qmmf11qmmf_kernelILi3ELi16ELi64ELi1ELi4ELi4ELb0ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii",
         symbol(.fp8g, 16, true, false, false),
diff --git a/zig/src/cuda/root.zig b/zig/src/cuda/root.zig
index 6c4e0991b6..f49c4de197 100644
--- a/zig/src/cuda/root.zig
+++ b/zig/src/cuda/root.zig
@@ -26,6 +26,7 @@ pub const aot = @import("aot.zig");
 pub const kernels = @import("kernels.zig");
 pub const qmmf = @import("qmmf.zig");
 pub const fp8 = @import("fp8.zig");
+pub const mxfp8 = @import("mxfp8.zig");
 pub const exlr8 = @import("exlr8.zig");
 pub const nvfp4 = @import("nvfp4.zig");
 pub const segments = @import("segments.zig");
@@ -58,6 +59,7 @@ test {
     _ = segments;
     _ = @import("qmmf.zig");
     _ = @import("fp8.zig");
+    _ = @import("mxfp8.zig");
     _ = carveout;
     _ = @import("nvfp4.zig");
     _ = @import("modelopt.zig");
diff --git a/zig/tests/cuda/contract_mxfp8.zig b/zig/tests/cuda/contract_mxfp8.zig
new file mode 100644
index 0000000000..278cb78c4c
--- /dev/null
+++ b/zig/tests/cuda/contract_mxfp8.zig
@@ -0,0 +1,115 @@
+//! MXFP8 projections through the common non-finite, repeat, shape and accumulator gates.
+const std = @import("std");
+const cuda = @import("cuda");
+const core = @import("core");
+const contract = core.contract;
+const c = contract.spec;
+const g = contract.gates;
+const Dense = contract.dense.Dense;
+const Rig = @import("contract_rig.zig").Rig;
+
+fn depth(shape: c.Shape) usize {
+    return shape.k / 16 + 24;
+}
+
+const served = [_]c.Shape{
+    .{ .rows = 1, .n = 70, .k = 1024 },
+    .{ .rows = 17, .n = 70, .k = 1024 },
+    .{ .rows = 33, .n = 70, .k = 1024 },
+    .{ .rows = 256, .n = 70, .k = 1024 },
+    .{ .rows = 16, .n = 1280, .k = 2560 },
+    .{ .rows = 16, .n = 2560, .k = 640 },
+    .{ .rows = 7, .n = 13952, .k = 2560 },
+};
+
+pub const Product = struct {
+    rig: Rig,
+    lane: *const cuda.qmmf.Lane,
+    clusters: bool,
+    d: ?Dense = null,
+    weight: cuda.qmmf.Weight = undefined,
+    x: u64 = 0,
+    out: u64 = 0,
+    part: u64 = 0,
+
+    pub fn deinit(self: *Product) void {
+        if (self.d) |d| d.deinit();
+        self.rig.deinit();
+    }
+
+    pub fn kernel(self: *Product) g.Kernel {
+        return .{ .spec = .{
+            .name = if (self.clusters) "cuda MXFP8 lane clustered slices" else "cuda MXFP8 lane explicit slice reduction",
+            .accum = .f32,
+            .depth = depth,
+            .out = .bf16,
+            .tiles = .{ .{ .size = 16 }, .{ .size = 64 }, .{ .size = 512, .step = 64, .min = 64 } },
+            .served = &served,
+        }, .ctx = self, .vt = &vt };
+    }
+
+    const vt = g.hooks(Product);
+
+    pub fn prepare(self: *Product, shape: c.Shape, fill: g.Fill, seed: u64) !bool {
+        if (self.d) |d| d.deinit();
+        self.d = null;
+        self.rig.reset();
+        self.part = 0;
+        if (shape.rows == 0 or shape.n == 0 or shape.k == 0 or shape.k % 64 != 0) return false;
+        const a = self.rig.gpu.gpa;
+        const d = try Dense.init(a, shape, 1);
+        self.d = d;
+        d.fillX(fill, seed, 16);
+        var prng = std.Random.DefaultPrng.init(seed ^ 0x8f);
+        const random = prng.random();
+        const n = shape.n;
+        const k = shape.k;
+        const codes = try a.alloc(u8, n * k);
+        defer a.free(codes);
+        const scales = try a.alloc(u8, n * k / 32);
+        defer a.free(scales);
+        for (scales) |*s| s.* = if (fill == .stress) 127 else 123 + random.uintLessThan(u8, 8);
+        for (0..n) |j| for (0..k) |i| {
+            codes[j * k + i] = if (fill == .stress) 0x38 + random.uintLessThan(u8, 24) else random.uintLessThan(u8, 0x7f);
+            if (fill != .stress and random.boolean()) codes[j * k + i] |= 0x80;
+        };
+        if (fill == .stress) for (0..n) |j| for (k / 2..k) |i| {
+            codes[j * k + i] = codes[j * k + i - k / 2] | 0x80;
+        };
+        for (0..n) |j| for (0..k) |i| {
+            const scale = std.math.ldexp(@as(f64, 1), @as(i32, scales[j * (k / 32) + i / 32]) - 127);
+            d.w[j * k + i] = core.quant.fp8.e4m3(codes[j * k + i]) * scale;
+        };
+        const npad = cuda.mxfp8.padded(n);
+        const tiled_codes = try a.alloc(u8, npad * k);
+        defer a.free(tiled_codes);
+        const packed_scales = try a.alloc(u8, npad * k / 32);
+        defer a.free(packed_scales);
+        cuda.mxfp8.packCodes(tiled_codes, codes, n, k);
+        cuda.mxfp8.packScales(packed_scales, scales, n, k);
+        self.weight = cuda.mxfp8.weight(try self.rig.input(tiled_codes), try self.rig.input(packed_scales), n, k);
+        self.x = try self.rig.input(std.mem.sliceAsBytes(d.x));
+        const bytes = self.lane.partBytes(shape.rows, self.weight);
+        if (bytes > 0) self.part = try self.rig.scratch(bytes);
+        self.out = try self.rig.output(shape.rows * n * 2);
+        return true;
+    }
+
+    pub fn run(self: *Product) !void {
+        const d = self.d.?;
+        try self.lane.matmul(self.rig.s, self.x, d.shape.k, d.shape.rows, self.weight, self.out, self.part);
+        try self.rig.s.synchronize();
+    }
+
+    pub fn count(self: *Product) usize {
+        return self.d.?.outputs();
+    }
+
+    pub fn reference(self: *Product, want: []f64, magnitude: []f64) !void {
+        try self.d.?.reference(want, magnitude);
+    }
+
+    pub fn simulate(self: *Product, accumulator: c.Accum, order: g.Order, out: []f64) !void {
+        try self.d.?.simulate(accumulator, order, out);
+    }
+};
diff --git a/zig/tests/cuda/contract_tests.zig b/zig/tests/cuda/contract_tests.zig
index e87a2e90e5..c89c814c6c 100644
--- a/zig/tests/cuda/contract_tests.zig
+++ b/zig/tests/cuda/contract_tests.zig
@@ -7,6 +7,7 @@ const check = @import("check.zig");
 const rows_k = @import("contract_rows.zig");
 const router_k = @import("contract_router.zig");
 const mm_k = @import("contract_mm.zig");
+const mx_k = @import("contract_mxfp8.zig");
 const rank_k = @import("contract_rank.zig");
 
 const contract = core.contract;
@@ -86,6 +87,12 @@ pub fn run(gpu: check.Gpu) !void {
     }
     defer for (&mm) |*m| m.deinit();
     for (&mm) |*m| try kernels.append(a, m.kernel());
+    var mx = [_]mx_k.Product{
+        .{ .rig = .{ .gpu = gpu, .s = s }, .lane = &lane, .clusters = true },
+        .{ .rig = .{ .gpu = gpu, .s = s }, .lane = &lane_reduce, .clusters = false },
+    };
+    defer for (&mx) |*m| m.deinit();
+    for (&mx) |*m| try kernels.append(a, m.kernel());
     var rank = rank_k.Rank{ .rig = .{ .gpu = gpu, .s = s }, .sum = &rank_sum };
     defer rank.deinit();
     try kernels.append(a, rank.kernel());
diff --git a/zig/tests/cuda/main.zig b/zig/tests/cuda/main.zig
index fe98682c86..89d267a5f5 100644
--- a/zig/tests/cuda/main.zig
+++ b/zig/tests/cuda/main.zig
@@ -17,6 +17,7 @@ const prompt_cuts = @import("prompt_cuts.zig");
 const chunk_costs = @import("chunk_costs.zig");
 const grouped_tests = @import("grouped_tests.zig");
 const fp8_tests = @import("fp8_tests.zig");
+const mxfp8_tests = @import("mxfp8_tests.zig");
 const fp8_experts_tests = @import("fp8_experts_tests.zig");
 const nvfp4_experts_tests = @import("nvfp4_experts_tests.zig");
 const quant_ref_tests = @import("quant_ref_tests.zig");
@@ -80,6 +81,7 @@ const usage =
     \\  fp8-lane <dir>            block-FP8 projections against the Python lane matmul's bytes (oracle/fp8_lane.py)
     \\  fp8-experts <dir>         grouped block-FP8 experts against recorded fixture bytes
     \\  nvfp4-experts <dir>       grouped NVFP4 experts against recorded fixture bytes
+    \\  mxfp8                     MXFP8 strided projections against FP64 and row-count invariance
     \\  quant-ref                 FP8 and NVFP4 projections and experts against float64 references (synthetic)
     \\  carveout [MiB] [card]     GB10 display memory: round trips and bandwidth (SKIP without the card)
     \\  nvfp4 <dir>               NVFP4 projections against recorded lane-matmul and prompt-GEMM fixture bytes
@@ -155,6 +157,7 @@ fn run(gpu: check.Gpu, cmd: []const u8, rest: []const [:0]const u8) !void {
     if (std.mem.eql(u8, cmd, "fp8-experts")) return fp8_experts_tests.experts(gpu, try arg(rest, 0));
     if (std.mem.eql(u8, cmd, "nvfp4")) return nvfp4_tests.run(gpu, try arg(rest, 0));
     if (std.mem.eql(u8, cmd, "nvfp4-experts")) return nvfp4_experts_tests.experts(gpu, try arg(rest, 0));
+    if (std.mem.eql(u8, cmd, "mxfp8")) return mxfp8_tests.run(gpu);
     if (std.mem.eql(u8, cmd, "quant-ref")) return quant_ref_tests.run(gpu);
     if (std.mem.eql(u8, cmd, "glm-ref")) return glm_ref_tests.run(gpu);
     if (std.mem.eql(u8, cmd, "modelopt")) return modelopt_tests.run(gpu);
diff --git a/zig/tests/cuda/mxfp8_tests.zig b/zig/tests/cuda/mxfp8_tests.zig
new file mode 100644
index 0000000000..5417d3d1fd
--- /dev/null
+++ b/zig/tests/cuda/mxfp8_tests.zig
@@ -0,0 +1,143 @@
+//! Native MXFP8: lossless packing, independent FP64 products, strided rows and row-count invariance.
+const std = @import("std");
+const cuda = @import("cuda");
+const core = @import("core");
+const check = @import("check.zig");
+
+fn bf(v: u16) f64 {
+    return @as(f32, @bitCast(@as(u32, v) << 16));
+}
+
+fn tobf(v: f32) u16 {
+    const u: u32 = @bitCast(v);
+    return @truncate((u + 0x7fff + ((u >> 16) & 1)) >> 16);
+}
+
+pub fn run(gpu: check.Gpu) !void {
+    try scaleRange(gpu);
+    const a = gpu.gpa;
+    const n = 70; // a partial tile and padded rows
+    const k = 1024; // split-K, both cluster and explicit reduction
+    const m = 256; // fused prompt reduction
+    const ldx = k + 64;
+    const npad = cuda.mxfp8.padded(n);
+    var prng = std.Random.DefaultPrng.init(0x8e80_5eed);
+    const rng = prng.random();
+    const codes = try a.alloc(u8, n * k);
+    defer a.free(codes);
+    for (codes) |*v| {
+        v.* = rng.int(u8);
+        if (v.* & 0x7f == 0x7f) v.* ^= 1;
+    }
+    const scales = try a.alloc(u8, n * (k / 32));
+    defer a.free(scales);
+    for (scales) |*v| v.* = 121 + rng.uintLessThan(u8, 8);
+    // A dedicated row makes UE8M0 byte 0 observable in a normal BF16 result.
+    @memset(codes[0..k], 0x38); // E4M3 1
+    @memset(scales[0 .. k / 32], 0);
+    const packed_codes = try a.alloc(u8, npad * k);
+    defer a.free(packed_codes);
+    cuda.mxfp8.packCodes(packed_codes, codes, n, k);
+    const packed_scales = try a.alloc(u8, npad * (k / 32));
+    defer a.free(packed_scales);
+    cuda.mxfp8.packScales(packed_scales, scales, n, k);
+    var dw = try cuda.DeviceBuffer.fromHost(gpu.d, packed_codes);
+    defer dw.free();
+    var ds = try cuda.DeviceBuffer.fromHost(gpu.d, packed_scales);
+    defer ds.free();
+    const w = cuda.mxfp8.weight(dw.ptr, ds.ptr, n, k);
+    const x = try a.alloc(u16, m * ldx);
+    defer a.free(x);
+    @memset(x, 0x7fc0); // stride padding is poison, not another activation
+    for (0..m) |row| {
+        for (x[row * ldx ..][0..k]) |*v| v.* = tobf(rng.float(f32) * 2);
+    }
+    var dx = try cuda.DeviceBuffer.fromHost(gpu.d, std.mem.sliceAsBytes(x));
+    defer dx.free();
+    var dy = try cuda.DeviceBuffer.alloc(gpu.d, m * n * 2);
+    defer dy.free();
+    var part = try cuda.DeviceBuffer.alloc(gpu.d, 8 * m * n * 4);
+    defer part.free();
+    var lane = try cuda.qmmf.Lane.load(gpu.d, (try gpu.ctx.capability()) / 10);
+    defer lane.unload();
+    var s = try cuda.Stream.init(gpu.d, false);
+    defer s.deinit();
+    var baseline: ?[]u8 = null;
+    defer if (baseline) |b| a.free(b);
+    for ([_]bool{ true, false }) |clusters| {
+        lane.clusters = clusters and (try gpu.ctx.capability()) >= 90;
+        const lin: cuda.qlinear.Linear = .{ .lane = &lane };
+        for ([_]usize{ 256, 1, 7, 17, 33 }) |rows| {
+            for ([_]bool{ false, true }) |prompt| {
+                try dy.fill8(0xff, s.handle);
+                const in: cuda.qlinear.Rows = .{ .x = dx.ptr, .ldx = ldx, .part = part.ptr };
+                if (prompt) try lin.prompt(s, .{ .mxfp8 = w }, in, dy.ptr, rows) else try lin.decode(s, .{ .mxfp8 = w }, in, dy.ptr, rows);
+                try s.synchronize();
+                const got = try a.alloc(u8, rows * n * 2);
+                defer a.free(got);
+                try dy.download(0, got);
+                if (baseline) |b| {
+                    try check.sameBytes("MXFP8 row-count/dispatch/cluster invariance", got, b[0..got.len]);
+                } else baseline = try a.dupe(u8, got);
+                for (0..rows) |row| for (0..n) |col| {
+                    var ref: f64 = 0;
+                    var magnitude: f64 = 0;
+                    for (0..k) |i| {
+                        const scale = std.math.ldexp(@as(f64, 1), @as(i32, scales[col * (k / 32) + i / 32]) - 127);
+                        const v = bf(x[row * ldx + i]) * core.quant.fp8.e4m3(codes[col * k + i]) * scale;
+                        ref += v;
+                        magnitude += @abs(v);
+                    }
+                    const actual = bf(std.mem.readInt(u16, got[(row * n + col) * 2 ..][0..2], .little));
+                    const tolerance = @abs(ref) / 128 + magnitude / 4096 + 0x1p-133;
+                    try check.expect(@abs(actual - ref) <= tolerance, "MXFP8 FP64 row {d}, column {d}: {e} versus {e}, tolerance {e}", .{ row, col, actual, ref, tolerance });
+                };
+            }
+        }
+    }
+    check.pass("MXFP8: strided rows 1,7,17,33,256 equal across dispatch and reduction paths; FP64 including UE8M0 byte zero", .{});
+}
+
+/// Every UE8M0 encoding, through the production matmul: each dot product is exactly one times its scale.
+fn scaleRange(gpu: check.Gpu) !void {
+    const a = gpu.gpa;
+    const n = 256;
+    const k = 64;
+    const codes = try a.alloc(u8, n * k);
+    defer a.free(codes);
+    @memset(codes, 0x38);
+    var scales: [n * 2]u8 = undefined;
+    for (0..n) |r| @memset(scales[r * 2 ..][0..2], @intCast(r));
+    const tiled_codes = try a.alloc(u8, n * k);
+    defer a.free(tiled_codes);
+    cuda.mxfp8.packCodes(tiled_codes, codes, n, k);
+    var tiles: [n * 2]u8 = undefined;
+    cuda.mxfp8.packScales(&tiles, &scales, n, k);
+    var dw = try cuda.DeviceBuffer.fromHost(gpu.d, tiled_codes);
+    defer dw.free();
+    var ds = try cuda.DeviceBuffer.fromHost(gpu.d, &tiles);
+    defer ds.free();
+    var x: [k]u16 = @splat(0x3c80); // 1/64: 64 equal terms sum to one
+    var dx = try cuda.DeviceBuffer.fromHost(gpu.d, std.mem.sliceAsBytes(&x));
+    defer dx.free();
+    var dy = try cuda.DeviceBuffer.alloc(gpu.d, n * 2);
+    defer dy.free();
+    var lane = try cuda.qmmf.Lane.load(gpu.d, (try gpu.ctx.capability()) / 10);
+    defer lane.unload();
+    var s = try cuda.Stream.init(gpu.d, false);
+    defer s.deinit();
+    try lane.matmul(s, dx.ptr, k, 1, cuda.mxfp8.weight(dw.ptr, ds.ptr, n, k), dy.ptr, 0);
+    try s.synchronize();
+    var got: [n * 2]u8 = undefined;
+    try dy.download(0, &got);
+    for (0..n) |r| {
+        const bits = std.mem.readInt(u16, got[r * 2 ..][0..2], .little);
+        if (r == 255) {
+            try check.expect(std.math.isNan(bf(bits)), "UE8M0 reserved byte produces NaN", .{});
+        } else {
+            const want: u16 = if (r == 0) 0x40 else @intCast(r << 7);
+            try check.expect(bits == want, "UE8M0 {d}: BF16 bits {x} versus {x}", .{ r, bits, want });
+        }
+    }
+    check.pass("MXFP8: all 256 UE8M0 encodings preserve the specified value, including subnormal and NaN", .{});
+}
diff --git a/zig/tests/cuda/quant_ref_host_tests.zig b/zig/tests/cuda/quant_ref_host_tests.zig
new file mode 100644
index 0000000000..dc13f4235e
--- /dev/null
+++ b/zig/tests/cuda/quant_ref_host_tests.zig
@@ -0,0 +1,83 @@
+//! CPU tests of Rig.close: only stream synchronization and readback are replaced; the comparison is the GPU gate's.
+const std = @import("std");
+const cuda = @import("cuda");
+const Rig = @import("quant_ref_tests.zig").RigT;
+
+const Readback = struct {
+    var syncs: usize = 0;
+    var copies: usize = 0;
+
+    fn sync(_: cuda.abi.Stream) callconv(.c) cuda.abi.Result {
+        syncs += 1;
+        return cuda.abi.success;
+    }
+
+    fn copy(dst: ?*anyopaque, src: cuda.abi.DevicePtr, n: usize) callconv(.c) cuda.abi.Result {
+        const out: [*]u8 = @ptrCast(dst.?);
+        const in: [*]const u8 = @ptrFromInt(src);
+        @memcpy(out[0..n], in[0..n]);
+        copies += 1;
+        return cuda.abi.success;
+    }
+};
+
+fn close(comptime T: type, got: []const T, want: []const f64, tol: []const f64) !void {
+    var d: cuda.Driver = undefined;
+    d.api.cuStreamSynchronize = Readback.sync;
+    d.api.cuMemcpyDtoH_v2 = Readback.copy;
+    var arena = std.heap.ArenaAllocator.init(std.testing.allocator);
+    defer arena.deinit();
+    var rig: Rig = .{
+        .gpu = .{ .d = &d, .ctx = undefined, .gpa = std.testing.allocator, .io = std.testing.io },
+        .a = arena.allocator(),
+        .s = .{ .d = &d, .handle = null },
+        .rng = undefined,
+    };
+    Readback.syncs = 0;
+    Readback.copies = 0;
+    const result = rig.close("CPU reference comparison", @intFromPtr(got.ptr), T == f32, want, tol);
+    try std.testing.expectEqual(@as(usize, 1), Readback.syncs);
+    try std.testing.expectEqual(@as(usize, 1), Readback.copies);
+    return result;
+}
+
+const zero = [_]f64{ 0, 0, 0 };
+const tolerance = [_]f64{ 1, 1, 1 };
+
+test "finite FP32 and BF16 outputs pass at the bound and fail beyond it" {
+    try close(f32, &.{ 0.5, 1, 0 }, &zero, &tolerance);
+    try close(u16, &.{ 0x3f00, 0x3f80, 0 }, &zero, &tolerance);
+    try std.testing.expectError(error.TestFailed, close(f32, &.{ 2, 0, 0 }, &zero, &tolerance));
+    try std.testing.expectError(error.TestFailed, close(u16, &.{ 0x4000, 0, 0 }, &zero, &tolerance));
+}
+
+test "FP32 NaN cannot disappear before a later finite result" {
+    for ([_]u32{ 0x7fc00000, 0xffffffff }) |bits| {
+        for (0..3) |i| {
+            var got = [_]f32{ 0, 0, 0 };
+            got[i] = @bitCast(bits);
+            try std.testing.expectError(error.TestFailed, close(f32, &got, &zero, &tolerance));
+        }
+    }
+}
+
+test "BF16 NaN cannot disappear before a later finite result" {
+    for ([_]u16{ 0x7fc0, 0xffff }) |bits| {
+        for (0..3) |i| {
+            var got = [_]u16{ 0, 0, 0 };
+            got[i] = bits;
+            try std.testing.expectError(error.TestFailed, close(u16, &got, &zero, &tolerance));
+        }
+    }
+}
+
+test "NaN cannot erase an earlier out-of-bound or infinite result" {
+    try std.testing.expectError(error.TestFailed, close(f32, &.{ 0, std.math.inf(f32), 0 }, &zero, &tolerance));
+    try std.testing.expectError(error.TestFailed, close(f32, &.{ 0, -std.math.inf(f32), 0 }, &zero, &tolerance));
+    try std.testing.expectError(error.TestFailed, close(u16, &.{ 0, 0x7f80, 0 }, &zero, &tolerance));
+    try std.testing.expectError(error.TestFailed, close(u16, &.{ 0, 0xff80, 0 }, &zero, &tolerance));
+    try std.testing.expectError(error.TestFailed, close(f32, &.{ 2, std.math.nan(f32), 0 }, &zero, &tolerance));
+    try std.testing.expectError(error.TestFailed, close(f32, &.{ std.math.inf(f32), std.math.nan(f32), 0 }, &zero, &tolerance));
+    try std.testing.expectError(error.TestFailed, close(u16, &.{ 0x4000, 0xffff, 0 }, &zero, &tolerance));
+    try std.testing.expectError(error.TestFailed, close(u16, &.{ 0x7f80, 0xffff, 0 }, &zero, &tolerance));
+}
diff --git a/zig/tests/cuda/quant_ref_tests.zig b/zig/tests/cuda/quant_ref_tests.zig
index 657cd388b7..a0f40b664c 100644
--- a/zig/tests/cuda/quant_ref_tests.zig
+++ b/zig/tests/cuda/quant_ref_tests.zig
@@ -58,6 +58,7 @@ const Rig = struct {
         for (want, tol, 0..) |w, t, i| {
             const g: f64 = if (f32_out) std.mem.bytesToValue(f32, raw[4 * i ..][0..4]) else bf(std.mem.bytesToValue(u16, raw[2 * i ..][0..2]));
             const q = @abs(g - w) / t;
+            try check.expect(std.math.isFinite(q), "{s}: non-finite error at element {d} (got {e}, reference {e}, tolerance {e})", .{ name, i, g, w, t });
             if (!(q <= most)) {
                 most = q;
                 at = i;