dsv41 Metal: router NaN guard in the MSL route kernels, and a dense Q4_K dequantizer; tests for both

Router: ds4_route and ds4_route_publish left an all-NaN gate row's pick at 0x7fffffff and read sc[0x7fffffff] (threadgroup memory) and
vram[0x7fffffff]; the test process dies with SIGBUS on the old kernel. ds4_route_fallback_pick takes the lowest expert not picked yet
(the CUDA/HIP guard's rule) and raises a flag in a buffer of the Metal host; router_nonfinite_take synchronises and reads it
(metal_host::router_bad_take), so ds4_main ends the request with the error it already has. The warning names the GPU on Metal; the
CUDA/HIP text is unchanged. Test: metal/tests/ds4_route_guard_parity.cpp (finite, all NaN, +-inf, half NaN, fewer than k finite).

Dense Q4_K: ds4_gguf.cu rows() sends a dense matrix of a type metal_host::dequant_rows does not know to iq::dequant_rows, whose CUDA
launch has no Metal dispatch and aborts the engine. ds4_dq_q4k (ggml's dequantize_row_q4_K, the Q5_K kernel without the fifth bit) is
added; any other dense type now fails with a plain message naming it. Test: metal/tests/ds4_dq_parity.cpp (Q4_K, Q5_K, Q6_K, Q8_0, BF16
against ref_quant, padded row stride).

metal/include/metal_compat/ds4_metal_host.h

@@ -100,6 +100,7 @@ bool mx_mv(const void* jobs, int nj, int max_n, size_t job_bytes, int fmt);
 bool q8_quant_rows(const float* x, size_t ldx, int rows, int k, int8_t* q, float* d);
 bool ring_to_ext(const void* ring, void* ext, long long p0, int n, int d, size_t row_bytes);
 bool route(const float* logits, const float* bias, int n, int k, int T, float route_scale, int* idx, float* w, const float* bias_vl, const void* vl);
+bool router_bad_take();   // a router kernel made up a pick (every score NaN) since the last call; clears it. Call after the stream is synchronised
 bool route_publish(const float* logits, const float* bias, int n, int k, int T, float route_scale, const void* vram, int* idx, float* w, void* plan_ptr, const float* bias_vl, unsigned vl);
 bool sample_greedy(const float* logits, int n_tokens, int n_vocab, const int* history, int history_len, int penalty_last_n, float penalty_repeat, float penalty_freq, float penalty_present, const float* logit_bias, int* out);
 bool sgemm_tn(const float* A, const float* B, float* C, int m, int nt, int k, int lda, int ldb, int ldc);

metal/src/engine/ds4_finish.cpp

@@ -118,8 +118,26 @@ bool ensure() {
     }
     return g_lib >= 0;
 }
+
+// One int the router kernels raise when a pick had to be made up (ds4_route_fallback_pick: every score NaN) - the Metal side of
+// ds4_kernels.cu's g_router_bad.  Shared storage, zeroed, kept for the process.
+mtl::Buffer g_route_bad;
+mtl::Buffer& route_bad() {
+    if (!g_route_bad.ptr) g_route_bad = g_cx->alloc(64);
+    return g_route_bad;
+}
 }  // namespace
 
+// Whether a router kernel made up a pick since the last call, and clears it.  The caller has synchronised the stream first (the
+// kernel's store is then visible); called with nothing dispatched yet it is false.
+bool router_bad_take() {
+    if (g_cx == nullptr || g_route_bad.ptr == nullptr) return false;
+    volatile int* f = (volatile int*) g_route_bad.ptr;
+    const bool bad = f[0] != 0;
+    if (bad) f[0] = 0;
+    return bad;
+}
+
 // the engine's Finish, as ds4_kernels.cuh declares it
 struct FinishIn {
     const float* w;
@@ -1347,7 +1365,8 @@ bool route(const float* logits, const float* bias, int n, int k, int T, float ro
     dsp.p = &pl;
     dsp.args = {mtl::Arg::B(bl, off(hl)), mtl::Arg::B(bb, off(hb)), mtl::Arg::V((unsigned) n), mtl::Arg::V((unsigned) k),
                 mtl::Arg::V(route_scale), mtl::Arg::B(bi, off(hi)), mtl::Arg::B(bw, off(hw)),
-                mtl::Arg::B(bv, hv != nullptr ? off(hv) : 0), mtl::Arg::V(has_vl ? 1u : 0u), mtl::Arg::B(bbv)};
+                mtl::Arg::B(bv, hv != nullptr ? off(hv) : 0), mtl::Arg::V(has_vl ? 1u : 0u), mtl::Arg::B(bbv),
+                mtl::Arg::B(route_bad())};
     dsp.grid[0] = (uint32_t) T;
     dsp.grid[1] = dsp.grid[2] = 1;
     dsp.tg[0] = 512;
@@ -2197,7 +2216,7 @@ bool route_publish(const float* logits, const float* bias, int n, int k, int T,
     d.args = {mtl::Arg::B(bl, off(hl)), mtl::Arg::B(bb, off(hb)), mtl::Arg::V((unsigned) n), mtl::Arg::V((unsigned) k),
               mtl::Arg::V(route_scale), mtl::Arg::B(bv, off(hv)), mtl::Arg::B(bi, off(hi)), mtl::Arg::B(bw, off(hw)),
               mtl::Arg::B(bp, off(hp)), mtl::Arg::B(bvl, vl_off), mtl::Arg::V((unsigned) vl),
-              mtl::Arg::V((unsigned) (bias_vl ? 1 : 0))};
+              mtl::Arg::V((unsigned) (bias_vl ? 1 : 0)), mtl::Arg::B(route_bad())};
     d.grid[0] = (uint32_t) T;
     d.grid[1] = d.grid[2] = 1;
     d.tg[0] = 512; d.tg[1] = d.tg[2] = 1;
@@ -2330,13 +2349,13 @@ bool rows_f16(const float* x, int ldx, const int* idx, int nrows, int ncols, uin
     return true;
 }
 
-// kind: 0 BF16 values, 8 Q8_0, 13 Q5_K, 14 Q6_K.  n rows of k values, ld the row stride (bytes; values for BF16), out half rows.
+// kind: 0 BF16 values, 8 Q8_0, 12 Q4_K, 13 Q5_K, 14 Q6_K.  n rows of k values, ld the row stride (bytes; values for BF16), out half rows.
 bool dequant_rows(int kind, const void* raw, size_t ld, int k, int n, uint16_t* out) {
     if (n <= 0 || k <= 0) return true;
     if (!ensure()) return false;
-    const char* fn = kind == 13 ? "ds4_dq_q5k" : kind == 14 ? "ds4_dq_q6k" : kind == 8 ? "ds4_dq_q8" : kind == 0 ? "ds4_dq_bf16" : nullptr;
+    const char* fn = kind == 12 ? "ds4_dq_q4k" : kind == 13 ? "ds4_dq_q5k" : kind == 14 ? "ds4_dq_q6k" : kind == 8 ? "ds4_dq_q8" : kind == 0 ? "ds4_dq_bf16" : nullptr;
     if (!fn) return false;
-    if (((kind == 13 || kind == 14) && k % 256 != 0) || (kind == 8 && k % 32 != 0)) return false;
+    if (((kind == 12 || kind == 13 || kind == 14) && k % 256 != 0) || (kind == 8 && k % 32 != 0)) return false;
     void* hw = strata::gpu::host_view((void*) raw);
     void* ho = strata::gpu::host_view(out);
     if (!hw || !ho) return false;
@@ -2347,7 +2366,7 @@ bool dequant_rows(int kind, const void* raw, size_t ld, int k, int n, uint16_t*
         return g_cx->wrap((void*) base, wrap_len(base, end - base));
     };
     auto off = [&](void* host) { return (size_t) ((uintptr_t) host - (((uintptr_t) host) & ~(uintptr_t) (page - 1))); };
-    const size_t wbytes = (size_t) (n - 1) * ld * (kind == 0 ? 2 : 1) + (kind == 0 ? (size_t) k * 2 : (kind == 13 ? (size_t) (k / 256) * 176 : kind == 14 ? (size_t) (k / 256) * 210 : (size_t) (k / 32) * 34));
+    const size_t wbytes = (size_t) (n - 1) * ld * (kind == 0 ? 2 : 1) + (kind == 0 ? (size_t) k * 2 : (kind == 12 ? (size_t) (k / 256) * 144 : kind == 13 ? (size_t) (k / 256) * 176 : kind == 14 ? (size_t) (k / 256) * 210 : (size_t) (k / 32) * 34));
     if (!std::getenv("DS4_NO_TOUCH")) {   // (the matrix may be file-backed pages not read yet: the GPU does not page them in, it stalls)
         volatile unsigned char sink = 0;
         for (size_t o = 0; o < wbytes; o += 4096) sink = sink + ((const volatile unsigned char*) hw)[o];
@@ -2358,7 +2377,7 @@ bool dequant_rows(int kind, const void* raw, size_t ld, int k, int n, uint16_t*
     if (!bw.handle || !bo.handle) return false;
     mtl::Pipeline pl;
     try { pl = g_cx->pipeline(g_lib, fn); } catch (...) { return false; }
-    size_t items = kind == 13 ? (size_t) n * (k / 256) * 8 : kind == 14 ? (size_t) n * (k / 256) * 64 : kind == 8 ? (size_t) n * (k / 32) : (size_t) n * k;
+    size_t items = kind == 12 || kind == 13 ? (size_t) n * (k / 256) * 8 : kind == 14 ? (size_t) n * (k / 256) * 64 : kind == 8 ? (size_t) n * (k / 32) : (size_t) n * k;
     mtl::Dispatch d;
     d.p = &pl;
     d.args = {mtl::Arg::B(bw, off(hw)), mtl::Arg::V((unsigned long) ld), mtl::Arg::V((unsigned) k), mtl::Arg::V((unsigned) n),

metal/src/engine/msl/ds4_init.metal

@@ -885,6 +885,23 @@ kernel void ds4_gemv_f32(device const float* W [[buffer(0)]], device const float
 }
 
 
+// ds4_kernels.cu route_fallback_pick - a pick that never came out: every score left is NaN (the gate's input went non-finite), so
+// nothing beat the start value and the index is still the 0x7fffffff it started as; the next line would read sc[0x7fffffff]
+// (threadgroup memory) and the kernels after it vram[0x7fffffff] / the plan.  The lowest expert not picked yet instead - a valid
+// index, so nothing reads outside the tables - and *bad raised for the host (router_bad_take, ds4_finish.cpp): the weights stay
+// NaN, the layer's output is garbage as it was, and the request hears of it.  Reached by a kernel's one thread, only when a pick is
+// missing, so a finite router pays one compare a pick.
+static inline int ds4_route_fallback_pick(threadgroup const int* pick, uint r, uint n, device int* bad) {
+    bad[0] = 1;
+    int ix = 0;
+    for (bool again = true; again;) {   // (the picks are few and unsorted: start over when one of them is the candidate)
+        again = false;
+        for (uint q = 0; q < r && !again; ++q)
+            if (pick[q] == ix) { ++ix; again = true; }
+    }
+    return ix < int(n) ? ix : 0;
+}
+
 // ds4_kernels.cu:1353  route_publish_kernel - the MoE router: sc = sqrt(softplus(logit)), bs = sc + bias, then k rounds of a
 // block-wide argmax (a strided max, a warp reduction, a 16-warp pass), the picks' sc summed into the weights, and each pick's
 // device address published to plan_ptr.  The prefetch requests (req) are NOT published here yet: that is the schedule, not the
@@ -895,6 +912,7 @@ kernel void ds4_route_publish(device const float* logits [[buffer(0)]], device c
                               device int* idx [[buffer(6)]], device float* w [[buffer(7)]],
                               device ulong* plan_ptr [[buffer(8)]], device const float* bias_vl [[buffer(9)]],
                               constant uint& vl [[buffer(10)]], constant uint& has_vl [[buffer(11)]],
+                              device int* bad [[buffer(12)]],
                               uint t [[threadgroup_position_in_grid]], uint tid [[thread_index_in_threadgroup]]) {
     threadgroup float sc[1024];
     threadgroup float bs[1024];
@@ -931,6 +949,7 @@ kernel void ds4_route_publish(device const float* logits [[buffer(0)]], device c
             for (uint q = 0; q < 16u; ++q) {
                 if (bv[q] > bb || (bv[q] == bb && bi[q] < ix)) { bb = bv[q]; ix = bi[q]; }
             }
+            if (ix >= int(n)) ix = ds4_route_fallback_pick(pick, r, n, bad);   // (a non-finite gate input)
             pick[r] = ix;
             bs[ix] = -3.402823466e+38f;
         }
@@ -2419,6 +2438,7 @@ kernel void ds4_route(device const float* logits [[buffer(0)]], device const flo
                       constant uint& n [[buffer(2)]], constant uint& k [[buffer(3)]], constant float& route_scale [[buffer(4)]],
                       device int* idx [[buffer(5)]], device float* w [[buffer(6)]], device const float* bias_vl [[buffer(7)]],
                       constant uint& has_vl [[buffer(8)]], device const uint* vlw [[buffer(9)]],
+                      device int* bad [[buffer(10)]],
                       uint tg [[threadgroup_position_in_grid]], uint tid [[thread_index_in_threadgroup]],
                       ushort lane [[thread_index_in_simdgroup]], ushort wp [[simdgroup_index_in_threadgroup]]) {
     threadgroup float sc[1024], bs[1024];
@@ -2465,6 +2485,7 @@ kernel void ds4_route(device const float* logits [[buffer(0)]], device const flo
                     b = bv[q];
                     ix = bi[q];
                 }
+            if (ix >= int(n)) ix = ds4_route_fallback_pick(pick, r, n, bad);   // (a non-finite gate input)
             pick[r] = ix;
             bs[ix] = -INFINITY;
         }
@@ -2706,6 +2727,33 @@ kernel void ds4_dq_q5k(device const uchar* W [[buffer(0)]], constant ulong& ld [
     }
 }
 
+// a thread a 32-value sub-block of a Q4_K matrix (144-byte blocks: d, dmin, 12 bytes of 6-bit scales and mins, 128 bytes of
+// nibbles; ggml's dequantize_row_q4_K: sub-block j takes the low nibbles of qs[32 * (j / 2) ..] for even j, the high ones for odd j,
+// d * scale * q - dmin * min) -> out[n][k] half.  The Q5_K kernel above without the fifth bit.
+kernel void ds4_dq_q4k(device const uchar* W [[buffer(0)]], constant ulong& ld [[buffer(1)]], constant uint& k [[buffer(2)]],
+                       constant uint& n [[buffer(3)]], device half* out [[buffer(4)]], uint gid [[thread_position_in_grid]]) {
+    const uint nb = k / 256u;
+    const ulong total = (ulong) n * nb * 8ul;
+    if (ulong(gid) >= total) return;
+    const uint ib = gid & 7u;
+    const uint blk = (gid >> 3) % nb;
+    const uint row = uint(ulong(gid) / (8ul * nb));
+    const device uchar* p = W + (ulong) row * ld + (ulong) blk * 144ul;
+    const float d = float(as_type<half>(*(const device ushort*) p));
+    const float dmin = float(as_type<half>(*(const device ushort*) (p + 2)));
+    int sc, mm;
+    ds4_kscale2(p + 4, int(ib), sc, mm);
+    const device uchar* qs = p + 16 + 32u * (ib >> 1);
+    const uint sh = 4u * (ib & 1u);
+    device packed_half4* o = (device packed_half4*) (out + (ulong) row * k + (ulong) blk * 256ul + ib * 32u);
+    const float ds = d * float(sc), dm = dmin * float(mm);
+    for (uint l = 0; l < 32u; l += 4u) {
+        const uint4 a = uint4(*(const device packed_uchar4*) (qs + l));
+        const uint4 q = (a >> sh) & 15u;
+        o[l >> 2] = packed_half4(half4(float4(ds) * float4(q) - float4(dm)));
+    }
+}
+
 // Q6_K (210-byte blocks: ql[128], qh[64], int8 scales[16], f16 d; the Unsloth sets' attention, shared-expert, head and engram_kv
 // matrices): a thread four values of a block (two 128-value halves x 32 lanes), ggml's dequantize_row_q6_K arithmetic
 // ((d * scale) * (q - 32), in this order) -> out[n][k] half

metal/tests/ds4_dq_parity.cpp

@@ -0,0 +1,67 @@
+// metal/tests/ds4_dq_parity.cpp - the DeepSeek engine's dense-matrix dequantizers to F16 rows (ds4_init.metal: ds4_dq_*; the host
+// side is dequant_rows in ds4_finish.cpp, reached from ds4_gguf.cu rows() for the prompt chunk's dense GEMM operands) against ggml's
+// dequantize_row_* (ref_quant.cpp).  Q4_K had none: a dense Q4_K matrix fell through to iq::dequant_rows, whose CUDA launch has no
+// Metal dispatch and ended the engine.  Rows are padded (ld above the row size) so the stride is tested too.
+//   cmake --build <build-metal> --target ds4_dq_parity && <build-metal>/ds4_dq_parity
+#include "mtl_rt.h"
+#include "ref_quant.h"
+
+#include <algorithm>
+#include <cmath>
+#include <cstdint>
+#include <cstdio>
+#include <cstring>
+#include <vector>
+
+int main() {
+    mtl::Context cx;
+    const int lib = cx.compile({"../src/engine/msl/ds4_init.metal"}, "", false);
+    struct Case { int type; const char* name; const char* fn; int k; };
+    const Case cases[] = {{ref::Q4_K, "Q4_K", "ds4_dq_q4k", 768}, {ref::Q4_K, "Q4_K", "ds4_dq_q4k", 256}, {ref::Q5_K, "Q5_K", "ds4_dq_q5k", 768},
+                          {ref::Q6_K, "Q6_K", "ds4_dq_q6k", 768}, {ref::Q8_0, "Q8_0", "ds4_dq_q8", 640}, {ref::BF16, "BF16", "ds4_dq_bf16", 640}};
+    const int n = 37;
+    int fails = 0;
+    for (const Case& c : cases) {
+        mtl::Pipeline pl = cx.pipeline(lib, c.fn);
+        const size_t row = ref::row_bytes(c.type, c.k);
+        const size_t ld = row + 64;   // bytes
+        std::vector<uint8_t> packed((size_t) n * row);
+        ref::random_blocks(c.type, packed.data(), n, c.k, 1234 + (uint64_t) c.type, 0.02f);
+        mtl::Buffer bw = cx.alloc(ld * (size_t) n), bo = cx.alloc((size_t) n * c.k * 2);
+        std::memset(bw.ptr, 0, ld * (size_t) n);
+        for (int r = 0; r < n; ++r) std::memcpy(bw.as<uint8_t>() + (size_t) r * ld, packed.data() + (size_t) r * row, row);
+        const bool bf = c.type == ref::BF16;
+        const unsigned long ldk = bf ? (unsigned long) (ld / 2) : (unsigned long) ld;   // BF16: values a row
+        size_t items;
+        switch (c.type) {
+            case ref::Q4_K: case ref::Q5_K: items = (size_t) n * (c.k / 256) * 8; break;
+            case ref::Q6_K: items = (size_t) n * (c.k / 256) * 64; break;
+            case ref::Q8_0: items = (size_t) n * (c.k / 32); break;
+            default: items = (size_t) n * c.k; break;
+        }
+        mtl::Dispatch d;
+        d.p = &pl;
+        d.args = {mtl::Arg::B(bw), mtl::Arg::V(ldk), mtl::Arg::V((uint32_t) c.k), mtl::Arg::V((uint32_t) n), mtl::Arg::B(bo)};
+        d.grid[0] = (uint32_t) ((items + 255) / 256);
+        d.tg[0] = 256;
+        cx.run({d});
+        std::vector<float> want((size_t) c.k);
+        double worst = 0;
+        int bad = 0;
+        for (int r = 0; r < n; ++r) {
+            ref::dequant(c.type, packed.data() + (size_t) r * row, want.data(), (size_t) c.k);
+            for (int i = 0; i < c.k; ++i) {
+                const float got = ref::f16_to_f32(bo.as<uint16_t>()[(size_t) r * c.k + i]);
+                const double err = std::fabs((double) got - (double) want[(size_t) i]);
+                const double tol = 1.5e-3 * std::fabs((double) want[(size_t) i]) + 1e-6;   // a half's rounding of the float result (2^-11 = 4.9e-4) and an fma's
+                worst = std::max(worst, err / (std::fabs((double) want[(size_t) i]) + 1e-6));
+                if (!(err <= tol)) ++bad;
+            }
+        }
+        std::printf("%-5s n=%d k=%d ld=%zu: %d of %d values outside half rounding of ggml's dequantize_row (worst rel %.2e)  %s\n", c.name, n,
+                    c.k, ld, bad, n * c.k, worst, bad == 0 ? "PASS" : "FAIL");
+        fails += bad != 0;
+    }
+    std::printf("%s\n", fails ? "ds4_dq_parity: FAILED" : "ds4_dq_parity: all ok");
+    return fails ? 1 : 0;
+}

metal/tests/ds4_route_guard_parity.cpp

@@ -0,0 +1,139 @@
+// metal/tests/ds4_route_guard_parity.cpp - the DeepSeek router's MSL kernels (ds4_init.metal: ds4_route, ds4_route_publish) on a
+// non-finite gate input.  The Metal twin of tests/dsv41/ds4_route_guard_test.cu: whatever the logits hold, the picks must be valid
+// expert indices (no read past the expert tables - before the guard an all-NaN row picked 0x7fffffff and read sc[0x7fffffff] /
+// vram[0x7fffffff]), the same as the host's top-k when the scores are finite, and the flag the host polls (ds4_finish.cpp's
+// router_bad_take) must be raised exactly when a pick had to be made up.
+//   cmake --build <build-metal> --target ds4_route_guard_parity && <build-metal>/ds4_route_guard_parity
+#include "mtl_rt.h"
+
+#include <algorithm>
+#include <cmath>
+#include <cstdint>
+#include <cstdio>
+#include <cstring>
+#include <utility>
+#include <vector>
+
+static int fails = 0;
+static void check(bool ok, const char* what) {
+    std::printf("%s %s\n", ok ? "ok  " : "FAIL", what);
+    if (!ok) ++fails;
+}
+
+int main() {
+    mtl::Context cx;
+    const int lib = cx.compile({"../src/engine/msl/ds4_init.metal"}, "", false);
+    mtl::Pipeline p_route = cx.pipeline(lib, "ds4_route"), p_pub = cx.pipeline(lib, "ds4_route_publish");
+
+    const int n = 384, k = 6, T = 16, kPicks = 8;
+    const float scale = 1.5f;
+    mtl::Buffer b_logits = cx.alloc((size_t) T * 1024 * 4), b_bias = cx.alloc(1024 * 4), b_idx = cx.alloc((size_t) T * kPicks * 4),
+                b_w = cx.alloc((size_t) T * kPicks * 4), b_vram = cx.alloc(1024 * 8), b_plan = cx.alloc((size_t) T * kPicks * 8),
+                b_vlw = cx.alloc(1024), b_bad = cx.alloc(64);
+    std::vector<float> bias(1024, 0.0f), logits((size_t) T * 1024, 0.0f);
+    for (int i = 0; i < n; ++i) bias[(size_t) i] = 0.01f * (float) ((i * 37) % 17);
+    std::vector<unsigned long long> vram(1024, 0ull);
+    for (int i = 0; i < n; i += 2) vram[(size_t) i] = 0x1000ull + (unsigned long long) i;
+    std::memcpy(b_bias.ptr, bias.data(), 1024 * 4);
+    std::memcpy(b_vram.ptr, vram.data(), 1024 * 8);
+
+    const float nan = std::nanf(""), inf = INFINITY;
+    const char* names[] = {"finite", "all NaN", "+inf everywhere", "-inf everywhere", "half NaN", "fewer than k finite"};
+    for (int mode = 0; mode < 6; ++mode) {
+        unsigned s = 12345u;
+        for (int t = 0; t < T; ++t)
+            for (int i = 0; i < 1024; ++i) {
+                s = s * 1664525u + 1013904223u;
+                float v = ((float) (s >> 8) / (float) (1 << 24) - 0.5f) * 8.0f;
+                if (i >= n) v = 0.0f;
+                if (mode == 1) v = nan;
+                if (mode == 2) v = inf;
+                if (mode == 3) v = -inf;
+                if (mode == 4 && (i & 1)) v = nan;
+                if (mode == 5 && i >= 3) v = nan;
+                logits[(size_t) t * 1024 + (size_t) i] = v;
+            }
+        std::memcpy(b_logits.ptr, logits.data(), logits.size() * 4);
+        for (int which = 0; which < 2; ++which) {
+            std::memset(b_idx.ptr, 0xff, (size_t) T * kPicks * 4);   // (-1: a pick the kernel never wrote is not a valid index)
+            std::memset(b_plan.ptr, 0, (size_t) T * kPicks * 8);
+            b_bad.as<int>()[0] = 0;
+            mtl::Dispatch d;
+            if (which == 0) {
+                d.p = &p_route;
+                d.args = {mtl::Arg::B(b_logits), mtl::Arg::B(b_bias), mtl::Arg::V((uint32_t) n), mtl::Arg::V((uint32_t) k),
+                          mtl::Arg::V(scale), mtl::Arg::B(b_idx), mtl::Arg::B(b_w), mtl::Arg::B(b_bias), mtl::Arg::V(0u),
+                          mtl::Arg::B(b_vlw), mtl::Arg::B(b_bad)};
+            } else {
+                d.p = &p_pub;
+                d.args = {mtl::Arg::B(b_logits), mtl::Arg::B(b_bias), mtl::Arg::V((uint32_t) n), mtl::Arg::V((uint32_t) k),
+                          mtl::Arg::V(scale), mtl::Arg::B(b_vram), mtl::Arg::B(b_idx), mtl::Arg::B(b_w), mtl::Arg::B(b_plan),
+                          mtl::Arg::B(b_bias), mtl::Arg::V(0u), mtl::Arg::V(0u), mtl::Arg::B(b_bad)};
+            }
+            d.grid[0] = (uint32_t) T;
+            d.tg[0] = 512;
+            cx.run({d});
+            const int* idx = b_idx.as<int>();
+            const float* w = b_w.as<float>();
+            const char* kn = which ? "route_publish" : "route";
+            char what[200];
+            bool valid = true, distinct = true, same = true;
+            for (int t = 0; t < T; ++t) {
+                for (int r = 0; r < k; ++r) {
+                    const int e = idx[(size_t) t * kPicks + r];
+                    valid = valid && e >= 0 && e < n;
+                    for (int q = 0; q < r; ++q) distinct = distinct && idx[(size_t) t * kPicks + q] != e;
+                }
+                if (mode == 0) {   // the host's top-k of the same scores
+                    std::vector<std::pair<float, int>> sc;
+                    for (int i = 0; i < n; ++i) {
+                        const float l = logits[(size_t) t * 1024 + (size_t) i];
+                        const float sp = l > 20.0f ? l : std::log1p(std::exp(l));
+                        sc.push_back({std::sqrt(sp) + bias[(size_t) i], i});
+                    }
+                    std::stable_sort(sc.begin(), sc.end(), [](const auto& a, const auto& b) { return a.first > b.first; });
+                    for (int r = 0; r < k; ++r) same = same && sc[(size_t) r].second == idx[(size_t) t * kPicks + r];
+                }
+            }
+            if (!valid || !distinct) {
+                std::printf("  token 0 picks:");
+                for (int r = 0; r < kPicks; ++r) std::printf(" %d", idx[r]);
+                std::printf("\n");
+            }
+            std::snprintf(what, sizeof what, "%s: %s picks are all in [0, %d)", names[mode], kn, n);
+            check(valid, what);
+            // (with fewer than k finite scores the picked entries' start value ties with the empty pick, so the router re-takes the lowest of
+            //  them: a repeat, harmless to the tables - and the weights are NaN there anyway - so only checked when k scores are usable)
+            if (mode != 1 && mode != 5) {
+                std::snprintf(what, sizeof what, "%s: %s picks are distinct", names[mode], kn);
+                check(distinct, what);
+            }
+            if (mode == 0) {
+                std::snprintf(what, sizeof what, "finite: %s picks equal the host's top-%d", kn, k);
+                check(same, what);
+                bool wsum = true;
+                for (int t = 0; t < T; ++t) {
+                    double sw = 0;
+                    for (int r = 0; r < k; ++r) sw += w[(size_t) t * kPicks + r];
+                    wsum = wsum && std::fabs(sw - scale) < 1e-4;
+                }
+                std::snprintf(what, sizeof what, "finite: %s weights sum to route_scale", kn);
+                check(wsum, what);
+            }
+            const bool flagged = b_bad.as<int>()[0] != 0;
+            std::snprintf(what, sizeof what, "%s: %s %s the non-finite flag", names[mode], kn, mode == 1 ? "raises" : "leaves clear");
+            check(flagged == (mode == 1), what);   // (fewer than k finite: the picked entries' start value re-wins, no stand-in is needed)
+            if (which == 1) {
+                const unsigned long long* pl = b_plan.as<unsigned long long>();
+                bool plan_ok = true;
+                for (int t = 0; t < T; ++t)
+                    for (int r = 0; r < k; ++r)
+                        plan_ok = plan_ok && pl[(size_t) t * kPicks + r] == vram[(size_t) idx[(size_t) t * kPicks + r]];
+                std::snprintf(what, sizeof what, "%s: route_publish's plan pointers are the table's entries of the picks", names[mode]);
+                check(plan_ok, what);
+            }
+        }
+    }
+    std::printf("%s\n", fails ? "ds4_route_guard_parity: FAILED" : "ds4_route_guard_parity: all ok");
+    return fails ? 1 : 0;
+}

src/dsv41/ds4_gguf.cu

@@ -1396,6 +1396,15 @@ void rows(const x3::Mat& m, uint16_t* dst, bool bf16, cudaStream_t s) {
         const int kind = m.fmt == kFmtGG ? m.K : m.fmt == kFmtQ8_0 ? 8 : m.fmt == kFmtBF16 ? 0 : -1;
         if (kind >= 0 && ::strata::dsv41::metal_host::dequant_rows(kind, m.raw, row_stride(m), m.k, m.n, dst)) return;
     }
+#endif
+#if defined(STRATA_METAL)
+    if (m.fmt == kFmtGG && m.raw != nullptr) {   // (what the Metal dequantizers above did not take - they cover Q4_K, Q5_K and Q6_K as F16 rows; the CUDA kernel below has no Metal dispatch and would abort the engine)
+        char what[192];
+        std::snprintf(what, sizeof what, "rows: a dense %s matrix as %s rows has no Metal dequantizer (Q4_K, Q5_K, Q6_K, Q8_0 and BF16 as F16 rows do)",
+                      iq::name(m.K), bf16 ? "BF16" : "F16");
+        fail(what);
+        return;
+    }
 #endif
     if (m.fmt == kFmtGG && m.raw != nullptr) {   // a K-quant / i-quant dense matrix (the IST-DAS sets' Q5_K): its own dequantizer
         if (!iq::dequant_rows(m.K, m.raw, row_stride(m), m.k, m.n, dst, bf16, s)) fail("rows: a low-bit type not covered");

src/dsv41/ds4_kernels.cu

@@ -2169,8 +2169,12 @@ void route_ahead(const float* logits, const float* bias, int n, int m, int T, Re
 }
 bool router_nonfinite_take(cudaStream_t s) {
 #if defined(STRATA_METAL)
-    (void) s;
-    return false;
+    // the MSL router kernels raise a flag in a buffer of the Metal host (ds4_route_fallback_pick); the stream's work is done first
+    if (cudaStreamSynchronize(s) != cudaSuccess) {
+        (void) cudaGetLastError();
+        return false;
+    }
+    return strata::dsv41::metal_host::router_bad_take();
 #else
     int v = 0;
     if (cudaStreamSynchronize(s) != cudaSuccess || cudaMemcpyFromSymbol(&v, g_router_bad, sizeof v) != cudaSuccess) {

src/dsv41/ds4_model.cu

@@ -8669,9 +8669,16 @@ struct Model::Impl {
             cudaSetDevice(d.id);
             if (router_nonfinite_take(d.st)) {
                 static std::atomic<bool> said{false};
-                if (!said.exchange(true))
-                    std::fprintf(stderr, "ds4: WARNING: the router met scores that are not numbers (NaN) on CUDA%d - the model's "
-                                         "activations went non-finite; its picks are stand-ins and the answer is not reliable\n", d.id);
+                if (!said.exchange(true)) {
+                    char where[32];   // (the same words on CUDA and HIP as before; the Metal engine has one GPU, named so)
+#if defined(STRATA_METAL)
+                    std::snprintf(where, sizeof where, "the GPU");
+#else
+                    std::snprintf(where, sizeof where, "CUDA%d", d.id);
+#endif
+                    std::fprintf(stderr, "ds4: WARNING: the router met scores that are not numbers (NaN) on %s - the model's "
+                                         "activations went non-finite; its picks are stand-ins and the answer is not reliable\n", where);
+                }
                 bad = true;
             }
         }