diff --git a/include/strata/core/verify.hpp b/include/strata/core/verify.hpp
index e6ff842..a13f7e4 100644
--- a/include/strata/core/verify.hpp
+++ b/include/strata/core/verify.hpp
@@ -547,6 +547,11 @@ private:
     float* h_ple_ = nullptr;     float* m_ple_ = nullptr;       // T * n_embd
     int32_t* h_out_ = nullptr;   int32_t* m_out_ = nullptr;     // T argmax ids
     uint32_t* h_nf_ = nullptr;   uint32_t* m_nf_ = nullptr;     // #879: kVerifyMaxT * kNonfiniteBlocks flags
+    // STRATA_DBG_NAN_WIN: per-layer non-finite flags scanned inside the window graph (device + host staging);
+    // entry l is layer l's input residual at pre(l), n_layers is the head's input, n_layers + 1 the mixed head input
+    uint32_t* dbg_lnf_ = nullptr;    uint32_t* dbg_lnf_h_ = nullptr;
+    int64_t dbg_lnf_stride_ = 0;     int64_t dbg_lnf_entries_ = 0;
+    float* dbg_lmax_ = nullptr;      float* dbg_lmax_h_ = nullptr;   // the same points' max |x| (one float an entry)
     int nf_row_ = -1;                     ///< nonfinite_row()
     int64_t nf_pos_ = -1;                 ///< nonfinite_pos()
     uint32_t nf_mask_ = 0;                ///< nonfinite_in(): bit t = row t
diff --git a/include/strata/kernels/verify_kernels.hpp b/include/strata/kernels/verify_kernels.hpp
index 45bc70e..1d1e658 100644
--- a/include/strata/kernels/verify_kernels.hpp
+++ b/include/strata/kernels/verify_kernels.hpp
@@ -175,6 +175,10 @@ void row_top_prob_split(const float* logits, int n_rows, int n_vocab, const int3
 /// be mapped host memory read after the window's own sync.  Graph-capturable; CUDA and HIP.
 inline constexpr int kNonfiniteBlocks = 16;
 void logits_nonfinite_rows(const float* logits, int n_rows, int n, uint32_t* flags, void* stream);
+/// STRATA_DBG_NAN_WIN: *out = max(*out, max |x[0..n)|) by atomic max on the ordered float bits.  `out` is one
+/// float per probe point, zeroed by the host before the window; a NaN shows as 0x7FFFFFFF-range bits (7FFFFFFF+
+/// ordered comparison still lands above every finite).  Graph-capturable; CUDA and HIP.
+void dbg_maxabs(const float* x, int64_t n, float* out, void* stream);
 /// True where row_top_prob_split runs as its own kernel (CUDA; AMD keeps the one-block row_top_prob, so it calls
 /// that).  STRATA_MULTI_BLOCK_ARGMAX=0: off.
 bool multi_block_head_ops();
diff --git a/src/core/verify.cpp b/src/core/verify.cpp
index e8f747f..a3958be 100644
--- a/src/core/verify.cpp
+++ b/src/core/verify.cpp
@@ -1091,6 +1091,8 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
 
     // ---- the window's inputs, from mapped staging.  With branches the steps and positions (first read by the first
     // layer's mixer) cross PCIe on side 1 beside the embedding and the first hyper-connection read.
+    if (dbg_lmax_ != nullptr)   // STRATA_DBG_NAN_WIN: the max probes start from 0 on every replay
+        cudaMemsetAsync(dbg_lmax_, 0, (size_t) dbg_lnf_entries_ * sizeof(float), cs);
     copy_i32_from_mapped(tok_, m_tok_, T, cs);
     bool inputs_pending = false;
     cudaStream_t in_s = cs;
@@ -1147,6 +1149,12 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
     // ---------------------------------------------------------------- pre(l, group): up to the ring
     auto pre = [&](int64_t l, int grp) -> bool {
         const int tb = tb_[grp], te = te_[grp], n = te - tb;
+        if (dbg_lnf_ != nullptr && grp == 0 && G == 1) {   // STRATA_DBG_NAN_WIN: layer l's input residual
+            strata::kernels::logits_nonfinite_rows(R_, (int) (T * HC), (int) N,
+                                                   dbg_lnf_ + (size_t) (l - lb_) * (size_t) dbg_lnf_stride_, cs);
+            if (dbg_lmax_ != nullptr)
+                strata::kernels::dbg_maxabs(R_, (int64_t) T * HC * N, dbg_lmax_ + (l - lb_), cs);
+        }
         stamp(l, 0, grp);
         const LayerView v(wt, l);
         const char* pfx[2] = {"hc_attn_", "hc_ffn_"};
@@ -1165,6 +1173,17 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
             if (grp == 0) {
                 if (ar_on()) wait_flag_ge(m_flag_, 1, cs);
                 copy_from_mapped(ple_, m_ple_, (int64_t) T * N, cs);
+                if (dbg_lnf_ != nullptr) {   // STRATA_DBG_NAN_WIN sub-scans: layer 1's GDN state slot, then the staged PLE row
+                    strata::kernels::logits_nonfinite_rows(ss.gdn_state + gdn_floats, (int) ((gdn_floats + N - 1) / N),
+                                                           (int) N, dbg_lnf_ + (size_t) (g.n_layers + 5) * (size_t) dbg_lnf_stride_, cs);
+                    strata::kernels::logits_nonfinite_rows(ple_, (int) T, (int) N,
+                                                           dbg_lnf_ + (size_t) (g.n_layers + 2) * (size_t) dbg_lnf_stride_, cs);
+                    if (dbg_lmax_ != nullptr) {
+                        strata::kernels::dbg_maxabs(ss.gdn_state + gdn_floats, (int64_t) gdn_floats,
+                                                    dbg_lmax_ + (g.n_layers + 5), cs);
+                        strata::kernels::dbg_maxabs(ple_, (int64_t) T * N, dbg_lmax_ + (g.n_layers + 2), cs);
+                    }
+                }
             }
             float* normalized = (float*) ((uint8_t*) ss.ple.scratch + ple_block_scratch_bytes());
             static const bool ple_batch_env = [] {
@@ -1258,6 +1277,12 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
                 if (!keep_rows) copy_from_mapped(hist_snap_ + (size_t) t * HS, hist, HS, cs);
             }
             pending = false;
+            if (dbg_lnf_ != nullptr && grp == 0) {   // STRATA_DBG_NAN_WIN sub-scan: R_ after the PLE block wrote it
+                strata::kernels::logits_nonfinite_rows(R_, (int) (T * HC), (int) N,
+                                                       dbg_lnf_ + (size_t) (g.n_layers + 3) * (size_t) dbg_lnf_stride_, cs);
+                if (dbg_lmax_ != nullptr)
+                    strata::kernels::dbg_maxabs(R_, (int64_t) T * HC * N, dbg_lmax_ + (g.n_layers + 3), cs);
+            }
         }
         auto gr_read_group = [&](int half, bool apply, float* inj_prev, float* inj_out) {
             FusedGrArgs fa[kFusedGrMaxT];
@@ -1986,6 +2011,12 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
                               !cvec().covers(g.n_layers - 1);
     auto post = [&](int64_t l, int grp) -> bool {
         const int tb = tb_[grp], te = te_[grp], n = te - tb;
+        if (dbg_lnf_ != nullptr && grp == 0 && G == 1 && l == 1) {   // STRATA_DBG_NAN_WIN sub-scan: R_ after layer 1's pre
+            strata::kernels::logits_nonfinite_rows(R_, (int) (T * HC), (int) N,
+                                                   dbg_lnf_ + (size_t) (g.n_layers + 4) * (size_t) dbg_lnf_stride_, cs);
+            if (dbg_lmax_ != nullptr)
+                strata::kernels::dbg_maxabs(R_, (int64_t) T * HC * N, dbg_lmax_ + (g.n_layers + 4), cs);
+        }
         const uint32_t ring = (uint32_t) ((l - lb_) * G + grp + 1);
         const int64_t cap = (int64_t) n * K, capx = (int64_t) max_t_ * K;
         int32_t* pl = plan_ + (size_t) grp * (size_t) (plan_i32_ + 16);
@@ -2146,6 +2177,12 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
 
     // ---- the head, T columns, and the argmax of each
     stamp(g.n_layers, 0, 0);
+    if (dbg_lnf_ != nullptr && G == 1) {   // STRATA_DBG_NAN_WIN: the head's input residual
+        strata::kernels::logits_nonfinite_rows(R_, (int) (T * HC), (int) N,
+                                               dbg_lnf_ + (size_t) g.n_layers * (size_t) dbg_lnf_stride_, cs);
+        if (dbg_lmax_ != nullptr)
+            strata::kernels::dbg_maxabs(R_, (int64_t) T * HC * N, dbg_lmax_ + g.n_layers, cs);
+    }
     {
         const WeightRef *hn = wt.find("output_hc_norm.weight"), *hd = wt.find("output_hc_down.weight"),
                         *hu = wt.find("output_hc_up.weight");
@@ -2211,6 +2248,12 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
             }
         }
         if (head_ != nullptr && head_->loaded()) {
+            if (dbg_lnf_ != nullptr && G == 1) {   // STRATA_DBG_NAN_WIN: the mixed head input, before the Q8_1 quantize
+                strata::kernels::logits_nonfinite_rows(head_mixed_, T, (int) N,
+                                                       dbg_lnf_ + (size_t) (g.n_layers + 1) * (size_t) dbg_lnf_stride_, cs);
+                if (dbg_lmax_ != nullptr)
+                    strata::kernels::dbg_maxabs(head_mixed_, (int64_t) T * N, dbg_lmax_ + (g.n_layers + 1), cs);
+            }
             try {
                 native_quantize_q8_1(head_mixed_, xq_, (int) N, T, cs);
                 if (g_mmvq_il() && strata::kernels::native_mmvq_il_supported(head_->type(), (int) T, (int) n_vocab_)) {
@@ -2290,6 +2333,28 @@ void Verifier::refresh_ar() {
 bool Verifier::capture(int T, std::string& err) {
     cudaGraphExec_t& exec_t = ar_off_ ? exec_nr_[T] : exec_[T];
     if (exec_t != nullptr) return true;
+    if (dbg_lnf_ == nullptr && g_ != nullptr && std::getenv("STRATA_DBG_NAN_WIN") != nullptr) {
+        // STRATA_DBG_NAN_WIN: one flags word block per layer entry, written by scans captured into the window graph.
+        // Stride 8192 words leaves room for the sub-scans (the GDN state slot's ~115 rows), entries 50+ below.
+        dbg_lnf_stride_ = 8192;
+        dbg_lnf_entries_ = g_->n_layers + 10;
+        const size_t bytes = (size_t) dbg_lnf_stride_ * (size_t) dbg_lnf_entries_ * sizeof(uint32_t);
+        if (cudaMalloc(&dbg_lnf_, bytes) == cudaSuccess && cudaMemset(dbg_lnf_, 0, bytes) == cudaSuccess) {
+            dbg_lnf_h_ = new uint32_t[(size_t) dbg_lnf_stride_ * (size_t) dbg_lnf_entries_]();
+            if (cudaMalloc(&dbg_lmax_, (size_t) dbg_lnf_entries_ * sizeof(float)) == cudaSuccess &&
+                cudaMemset(dbg_lmax_, 0, (size_t) dbg_lnf_entries_ * sizeof(float)) == cudaSuccess)
+                dbg_lmax_h_ = new float[(size_t) dbg_lnf_entries_]();
+            else {
+                dbg_lmax_ = nullptr;
+                (void) cudaGetLastError();
+            }
+            std::fprintf(stderr, "strata verify: STRATA_DBG_NAN_WIN on: %lld scan entries (per-window layer inputs)\n",
+                         (long long) dbg_lnf_entries_);
+        } else {
+            dbg_lnf_ = nullptr;
+            (void) cudaGetLastError();
+        }
+    }
     {   // said before the capture: a process that exits inside it (#1275: Windows, 313 MiB free) leaves this line as the trace
         size_t free_b = 0, total_b = 0;
         if (cudaMemGetInfo(&free_b, &total_b) == cudaSuccess)
@@ -3005,7 +3070,10 @@ bool Verifier::run(int T, const int32_t* tokens, int64_t pos0, PoolMultiFn pool,
     }
     for (int t = 0; t < T; ++t) out[t] = ((volatile int32_t*) h_out_)[t];
     note_nonfinite(T, nullptr);
-    if (static const bool dbg = std::getenv("STRATA_DBG_NAN") != nullptr; dbg) {   // debug: the first non-finite head
+    if (static const bool dbg = std::getenv("STRATA_DBG_NAN") != nullptr; dbg ||
+        (std::getenv("STRATA_DBG_NAN_END") != nullptr && nf_row_ >= 0) ||
+        (dbg_lnf_ != nullptr && nf_row_ >= 0)) {
+        // With the lightweight hand-off probe, only copy logits after the normal guard has fired.
         static bool reported = false;
         if (!reported) {
             std::vector<float> h((size_t) T * (size_t) n_vocab_);
@@ -3019,6 +3087,54 @@ bool Verifier::run(int T, const int32_t* tokens, int64_t pos0, PoolMultiFn pool,
                                          "(token out %d)\n", (long long) pos0, t, (long long) bad, (long long) n_vocab_, out[t]);
                 }
             }
+            if (nf_row_ >= 0 && dbg_lnf_ != nullptr && !reported) reported = true;   // the layer scan is the report
+            if (nf_row_ >= 0 && dbg_lnf_ != nullptr) {
+                const int64_t valid = (int64_t) T * g_->hc * strata::kernels::kNonfiniteBlocks;   // this T's rows
+                if (cudaMemcpy(dbg_lnf_h_, dbg_lnf_, (size_t) dbg_lnf_stride_ * (size_t) dbg_lnf_entries_ * 4,
+                               cudaMemcpyDeviceToHost) == cudaSuccess) {
+                    int64_t first = -1;
+                    std::string line;
+                    char b[96];
+                    for (int64_t e = 0; e < dbg_lnf_entries_; ++e) {
+                        const int64_t words = e < g_->n_layers + 2 ? std::min<int64_t>(valid, dbg_lnf_stride_)
+                                                                   : dbg_lnf_stride_;   // sub-scans: their own row counts
+                        int64_t bad = 0;
+                        for (int64_t w = 0; w < words; ++w)
+                            bad += dbg_lnf_h_[(size_t) e * (size_t) dbg_lnf_stride_ + (size_t) w] != 0;
+                        if (bad > 0) {
+                            if (first < 0) first = e;
+                            const char* nm = e == g_->n_layers + 2 ? "PLErow" : e == g_->n_layers + 3 ? "RafterPLE"
+                                         : e == g_->n_layers + 4 ? "RafterPre1" : e == g_->n_layers + 5 ? "GDNstate1" : "?";
+                            std::snprintf(b, sizeof b, " %s%lld:%lldw", e < g_->n_layers ? "L" : e < g_->n_layers + 2 ? "H" : "",
+                                          e < g_->n_layers ? (long long) e
+                                          : e < g_->n_layers + 2 ? e - g_->n_layers
+                                          : (long long) 0, (long long) bad);
+                            if (e >= g_->n_layers + 2) { std::snprintf(b, sizeof b, " %s:%lldw", nm, (long long) bad); }
+                            line += b;
+                        }
+                    }
+                    std::fprintf(stderr, "strata dbg: STRATA_DBG_NAN_WIN layer scan at position %lld (T=%d): entry = the "
+                                         "layer's input residual; L%lld = layer %lld's input, H0 = the head's input, "
+                                         "H1 = the mixed head input. First non-finite entry: %lld. Bad words:%s\n",
+                                 (long long) pos0, T, (long long) first,
+                                 first >= 0 && first < g_->n_layers ? (long long) first : -1LL, (long long) first,
+                                 line.empty() ? " none (every layer input finite: the NaN starts inside a layer's own "
+                                                "compute or the head)" : line.c_str());
+                    if (dbg_lmax_ != nullptr && dbg_lmax_h_ != nullptr &&
+                        cudaMemcpy(dbg_lmax_h_, dbg_lmax_, (size_t) dbg_lnf_entries_ * sizeof(float),
+                                   cudaMemcpyDeviceToHost) == cudaSuccess) {
+                        std::string mx;
+                        char c[64];
+                        for (int64_t e = 0; e < dbg_lnf_entries_; ++e) {
+                            if (e == g_->n_layers) mx += " | head:";
+                            std::snprintf(c, sizeof c, " %lld=%g", (long long) e, dbg_lmax_h_[(size_t) e]);
+                            mx += c;
+                        }
+                        std::fprintf(stderr, "strata dbg: STRATA_DBG_NAN_WIN max|x| per entry (50=PLErow 51=RafterPLE "
+                                             "52=RafterPre1 53=GDNstate1):%s\n", mx.c_str());
+                    }
+                }
+            }
         }
     }
     VDBG("window done\n");
diff --git a/src/kernels/cuda/verify_kernels.cu b/src/kernels/cuda/verify_kernels.cu
index a451d9b..4e09bd6 100644
--- a/src/kernels/cuda/verify_kernels.cu
+++ b/src/kernels/cuda/verify_kernels.cu
@@ -943,6 +943,37 @@ void logits_nonfinite_rows(const float* logits, int n_rows, int n, uint32_t* fla
     check("logits_nonfinite_rows");
 }
 
+namespace {
+// STRATA_DBG_NAN_WIN: max |x| over n floats by atomic max on the ordered bits; fabsf makes a NaN positive, whose
+// bits sit above +inf's, so the winner reads back as nan/inf when it is one.
+__global__ void dbg_maxabs_kernel(const float* __restrict__ x, long long n, int* __restrict__ out) {
+    __shared__ int smax[256];
+    const long long i0 = (long long) blockIdx.x * blockDim.x + threadIdx.x;
+    const long long stride = (long long) gridDim.x * blockDim.x;
+    int m = 0;
+    for (long long i = i0; i < n; i += stride) {
+        const int v = (int) __float_as_uint(fabsf(x[i]));
+        if (v > m) m = v;
+    }
+    smax[threadIdx.x] = m;
+    __syncthreads();
+    for (int s = (int) blockDim.x / 2; s > 0; s >>= 1) {
+        if ((int) threadIdx.x < s && smax[threadIdx.x + s] > smax[threadIdx.x]) smax[threadIdx.x] = smax[threadIdx.x + s];
+        __syncthreads();
+    }
+    if (threadIdx.x == 0) atomicMax(out, smax[0]);
+}
+}  // namespace
+
+void dbg_maxabs(const float* x, int64_t n, float* out, void* stream) {
+    if (x == nullptr || out == nullptr || n <= 0) return;
+    const int threads = 256;
+    const long long want = ((long long) n + threads - 1) / threads;
+    const long long cap = want < 4096 ? want : 4096;
+    dbg_maxabs_kernel<<<(unsigned) cap, threads, 0, (cudaStream_t) stream>>>(x, (long long) n, (int*) out);
+    check("dbg_maxabs");
+}
+
 bool multi_block_head_ops() {
 #if defined(__HIPCC__)
     return false;
diff --git a/src/prefill/prefill.cpp b/src/prefill/prefill.cpp
index f00f1cc..0bb2f62 100644
--- a/src/prefill/prefill.cpp
+++ b/src/prefill/prefill.cpp
@@ -3039,6 +3039,21 @@ bool Prefill::run_impl(const int64_t* tokens, int64_t n, int64_t pos0, std::stri
                 std::fprintf(stderr, "strata pf-step: chunk from token %lld, layer %lld: %s took %.0f ms (%s / %s)\n",
                              (long long) p0, (long long) layer, what, ms, cudaGetErrorString(a), cudaGetErrorString(b));
         };
+        auto nan_stage_probe = [&](const char* label, const float* ptr, int64_t rows, int64_t cols, int64_t start) {
+            if (std::getenv("STRATA_DBG_NAN_WIN") == nullptr) return;
+            cudaStreamSynchronize(m.cs);
+            std::vector<float> host((size_t) rows * cols);
+            if (cudaMemcpy(host.data(), ptr, host.size() * sizeof(float), cudaMemcpyDeviceToHost) != cudaSuccess) return;
+            int64_t bad = 0, first = -1;
+            double mx = 0;
+            for (size_t j = 0; j < host.size(); ++j) {
+                if (!std::isfinite(host[j])) { if (first < 0) first = (int64_t) j; ++bad; }
+                else mx = std::max(mx, (double) std::fabs(host[j]));
+            }
+            std::fprintf(stderr, "strata dbg stage: %s start=%lld T=%lld: %lld non-finite (first token %lld col %lld), max finite |x| %.3g\n",
+                         label, (long long) start, (long long) rows, (long long) bad,
+                         (long long) (first < 0 ? -1 : start + first / cols), (long long) (first < 0 ? -1 : first % cols), mx);
+        };
         for (int64_t l = LB; l < LE; ++l) {
             core::progress_beat();   // the serve watchdog: a prompt chunk of 8192 tokens is still moving
             if (l > LB) pf_step("reading the prompt (batched, step sync): the experts and the rest of layer", l - 1);
@@ -3049,6 +3064,10 @@ bool Prefill::run_impl(const int64_t* tokens, int64_t n, int64_t pos0, std::stri
                 if (!ple_land()) return false;
             }
             // ---- the PLE block at layer 1, token by token (its conv reads the previous tokens' rows)
+            if (l == 1 && ple_on) {
+                nan_stage_probe("pre-PLE residual", m.R, T, D, p0);
+                nan_stage_probe("PLE embedding", m.ple_emb, T, N, p0);
+            }
             if (l == 1 && ple_on && ple_batch) {
                 // the whole chunk at once, in sub-batches carved from the idle scratch region: the key and value
                 // projections as GEMMs (a token at a time they re-read ~52 MB of BF16 key per token on the IQ
@@ -3082,10 +3101,15 @@ bool Prefill::run_impl(const int64_t* tokens, int64_t n, int64_t pos0, std::stri
                     }
                     m.gemm.bf16(e16, pw.value_bf16, val, nb, N, N);
                     if (e16_lo) m.gemm.bf16(e16_lo, pw.value_bf16, val, nb, N, N, 0, 1.0f);
+                    nan_stage_probe("PLE projected key", key, nb, HD, p0 + s0);
+                    nan_stage_probe("PLE projected value", val, nb, N, p0 + s0);
                     try {
                         strata::kernels::native_ple_postops_batch(key, m.R + s0 * D, val, ss.ple.hist, pw, qn, gated,
                                                                   gate, (int) nb, m.cs);
                     } catch (const std::exception& e) { err = std::string("prefill PLE: ") + e.what(); return false; }
+                    nan_stage_probe("PLE normalized query", qn, nb, HD, p0 + s0);
+                    nan_stage_probe("PLE gated value", gated, nb, HD, p0 + s0);
+                    nan_stage_probe("PLE gate", gate, nb, HC, p0 + s0);
                 }
                 stats_.ms_ple += ms_since(tp);
             } else if (l == 1 && ple_on) {
@@ -3103,6 +3127,7 @@ bool Prefill::run_impl(const int64_t* tokens, int64_t n, int64_t pos0, std::stri
                 }
                 stats_.ms_ple += ms_since(tp);
             }
+            if (l == 1 && ple_on) nan_stage_probe("post-PLE residual", m.R, T, D, p0);
             for (int half = 0; half < 2; ++half) {
                 // ---- the hyper-connection read of this half
                 const char* pre = half == 0 ? "hc_attn_" : "hc_ffn_";
@@ -3182,6 +3207,23 @@ bool Prefill::run_impl(const int64_t* tokens, int64_t n, int64_t pos0, std::stri
                              m.mixed_bf_lo);
                 }
 
+                if (l == 1 && std::getenv("STRATA_DBG_NAN_WIN") != nullptr) {
+                    cudaStreamSynchronize(m.cs);
+                    std::vector<float> input((size_t) T * N);
+                    if (cudaMemcpy(input.data(), m.mixed, input.size() * sizeof(float), cudaMemcpyDeviceToHost) == cudaSuccess) {
+                        int64_t bad = 0, first = -1;
+                        double mx = 0;
+                        for (size_t j = 0; j < input.size(); ++j) {
+                            if (!std::isfinite(input[j])) { if (first < 0) first = (int64_t) j; ++bad; }
+                            else mx = std::max(mx, (double) std::fabs(input[j]));
+                        }
+                        std::fprintf(stderr, "strata dbg: layer1 %s input p0=%lld T=%lld: %lld non-finite "
+                                             "(first token +%lld col %lld), max finite |x| %.3g\n",
+                                     half == 0 ? "GDN" : "MoE", (long long) p0, (long long) T,
+                                     (long long) bad, (long long) (first < 0 ? -1 : first / N),
+                                     (long long) (first < 0 ? -1 : first % N), mx);
+                    }
+                }
                 if (wd_q8.data) {   // the HC GEMMs that read the scratch are enqueued: the next user waits for them
                     core::NativeDense::hc_bf16_release(m.cs, 0, 0);
                     core::NativeDense::hc_bf16_release(m.cs, 1, 0);
@@ -3208,6 +3250,40 @@ bool Prefill::run_impl(const int64_t* tokens, int64_t n, int64_t pos0, std::stri
                     const int64_t ld_y = pf_pad() && T >= std::max<int64_t>(pf_switch_min_t(), 64) ? ZV + ZV_PAD : 0;
                     gdn_recurrence(state, m.hbuf, m.gate, m.beta, m.z, (const float*) wnm->data, EPS, m.y, m.y_h, T, m.cs,
                                    ld_y);
+                    if (std::getenv("STRATA_DBG_NAN_WIN") != nullptr && gdn_index == 1) {
+                        // The first slot the #879 poison shows in: layer 1's GDN inputs and outputs, token by token,
+                        // so the non-finite value is pinned to one tensor and one prompt position.
+                        cudaStreamSynchronize(m.cs);
+                        static std::vector<float> dbgb;
+                        auto scan = [&](const char* what, const float* d, int64_t rows, int64_t cols) {
+                            if (rows <= 0 || cols <= 0) return;
+                            const size_t n = (size_t) rows * (size_t) cols;
+                            if (dbgb.size() < n) dbgb.resize(n);
+                            if (cudaMemcpy(dbgb.data(), d, n * 4, cudaMemcpyDeviceToHost) != cudaSuccess) return;
+                            int64_t bad = 0, first_row = -1, first_col = -1;
+                            double mx = 0;
+                            for (int64_t r = 0; r < rows; ++r)
+                                for (int64_t c = 0; c < cols; ++c) {
+                                    const float x = dbgb[(size_t) r * (size_t) cols + (size_t) c];
+                                    if (!std::isfinite(x)) {
+                                        ++bad;
+                                        if (first_row < 0) { first_row = r; first_col = c; }
+                                    } else mx = std::max(mx, (double) std::fabs((double) x));
+                                }
+                            if (bad > 0 || mx > 65504.0 || p0 == 0)
+                                std::fprintf(stderr, "strata dbg: GDN1 %s p0=%lld T=%lld: %lld non-finite (first at "
+                                                     "token +%lld col %lld), max finite |x| %.3g\n", what,
+                                             (long long) p0, (long long) T, (long long) bad, (long long) first_row,
+                                             (long long) first_col, mx);
+                        };
+                        scan("qkv", m.qkv, T, (int64_t) g.ssm_conv_channels);
+                        scan("hbuf", m.hbuf, T, (int64_t) g.ssm_conv_channels);
+                        scan("gate", m.gate, T, (int64_t) g.ssm_v_heads);
+                        scan("beta", m.beta, T, (int64_t) g.ssm_v_heads);
+                        scan("z", m.z, T, (int64_t) g.ssm_value_dim);
+                        scan("y", m.y, T, (int64_t) g.ssm_v_heads * (int64_t) g.ssm_state_size);
+                        scan("state", state, g.ssm_state_size, (int64_t) g.ssm_state_size * (int64_t) g.ssm_v_heads);
+                    }
                     pt.mark(kPfGdnOut, cs);
                     if (!native_proj(m.gemm, wo, m.y_h, m.bo, T, v.name("ssm_out.weight"), err, 0, ld_y)) return false;
                     ++gdn_index;
@@ -4919,7 +4995,9 @@ bool Prefill::run_impl(const int64_t* tokens, int64_t n, int64_t pos0, std::stri
     // stage can return while later GPUs are still processing the previous chunk.
     ss.ple_prev[0] = prev[0];
     ss.ple_prev[1] = prev[1];
-    if (std::getenv("STRATA_DBG_NAN") != nullptr) {   // debug: the state the prompt leaves for the token path
+    if (std::getenv("STRATA_DBG_NAN") != nullptr || std::getenv("STRATA_DBG_NAN_END") != nullptr ||
+        std::getenv("STRATA_DBG_NAN_WIN") != nullptr) {
+        // Inspect the hand-off without the expensive per-layer MoE scans enabled by STRATA_DBG_NAN.
         cudaStreamSynchronize(m.cs);
         auto bad = [&](const float* d, int64_t n) {
             std::vector<float> h((size_t) n);
@@ -4936,6 +5014,25 @@ bool Prefill::run_impl(const int64_t* tokens, int64_t n, int64_t pos0, std::stri
         std::fprintf(stderr, "; GDN state 0");
         bad(ss.gdn_state, 64 * 1024);
         std::fprintf(stderr, "\n");
+        // STRATA_DBG_NAN_WIN: every GDN slot (state + conv), as each chunk leaves it. A state can go non-finite
+        // while the outputs stay finite; the poison then surfaces in the next window's recurrence (#879).
+        std::vector<float> slot((size_t) gdn_floats);
+        for (int64_t i = 0; i < ss.gdn_alloc; ++i) {
+            if (cudaMemcpy(slot.data(), ss.gdn_state + (size_t) i * gdn_floats, slot.size() * 4,
+                           cudaMemcpyDeviceToHost) != cudaSuccess)
+                break;
+            int64_t c = 0, c_at = -1;
+            double mx = 0;
+            for (int64_t v = 0; v < gdn_floats; ++v) {
+                const float x = slot[(size_t) v];
+                if (!std::isfinite(x)) { if (c == 0) c_at = v; ++c; }
+                else mx = std::max(mx, (double) std::fabs((double) x));
+            }
+            if (c > 0 || mx > 65504.0)
+                std::fprintf(stderr, "strata dbg: GDN slot %lld after %lld tokens: %lld non-finite (first at %lld), "
+                                     "max finite |x| %.3g\n", (long long) i, (long long) n, (long long) c,
+                             (long long) c_at, mx);
+        }
     }
     if (cudaStreamSynchronize(m.cs) != cudaSuccess) {
         err = std::string("prefill: ") + cudaGetErrorString(cudaGetLastError());