#606: clamp the three unclamped q8_1 scale/sum emit sites via q8_1_ds

Three kernels still emit the (d, sum) pair of a q8_1 activation block
with a raw make_half2(d, sum), bypassing the finite helper the rest of
the tree uses (q8_1_finite.hpp, from 19de7f4 / f6601bf):

  fused_gr.cu   gr_q8_tail                 - the fused gate-reduce tail (QFUSE)
  verify_kernels.cu  gdn_q8_1_store        - the batched GDN verify store (QFUSE)
  iq_kernels.cu  s26_swiglu_q8_1_kernel   - the S26 swiglu path

A block whose amax overflows (or that reads a non-finite input) emits
d = inf/NaN; fp16 makes the pair NaN (or the pathological (0,0)). Nothing
downstream tests for it: a NaN scale poisons the fp16 dot of the quantized
activation and every sum downstream of it, and a fully non-finite logit
row ends at the sampler's all-NaN fallback - vocabulary id 0, the '!'
face of the degeneration reported in #606. Same defect class already
fixed at the other emit sites; these three were missed.

q8_1_ds() clamps both halves into the fp16-finite range (finite d,
non-negative finite sum - the sum is a mean of |x|), producing for the
GPU paths the same bits the CPU pool and the clamped GPU sites produce
for the same block. A clamp at the producer, content-blind: nothing
downstream of it is model-text-aware.

Criteria and the wider carried-state RCA: docs/Q8_1_FINITE_DS.md and
docs/DEGENERATION_RCA.md.

src/kernels/cuda/fused_gr.cu

@@ -1,5 +1,6 @@
 // src/kernels/cuda/fused_gr.cu - see include/strata/kernels/fused_gr.hpp.
 #include "strata/core/emulate.hpp"
+#include "strata/kernels/q8_1_finite.hpp"   // #606: q8_1_ds
 #include "strata/kernels/fused_gr.hpp"
 #include "strata/kernels/bf16_bits.hpp"
 #include "strata/kernels/verify_kernels.hpp"
@@ -199,7 +200,7 @@ __device__ __forceinline__ void gr_q8_tail(const GrMulti& m, int d0) {
     const int8_t q = amax == 0.0f ? 0 : roundf(xi / d);
     GrQ81* y = reinterpret_cast<GrQ81*>(m.a[warp].q8_mixed) + c0 / 32;
     y->qs[lane] = q;
-    if (lane == 0) y->ds = make_half2(d, sum);
+    if (lane == 0) y->ds = q8_1_ds(d, sum);   // #606: clamped scale/sum - an unclamped pair NaN-poisons the dot path
 }
 // Step 1 of `gr_down_kernel`, one block per token, same threads and reduction order: rs[t] and xn[t] to global.
 __global__ void __launch_bounds__(THREADS) gr_norm_multi_kernel(GrMulti m) {

src/kernels/cuda/iq_kernels.cu

@@ -2327,7 +2327,7 @@ __global__ void s26_swiglu_q8_1_kernel(const float* __restrict__ gate, const flo
     const int8_t q = amax == 0.0f ? 0 : roundf(xi / d);
     const long long ib = i / 32, iqs = i % 32;
     y[ib].qs[iqs] = q;
-    if (iqs == 0) y[ib].ds = make_half2(d, sum);
+    if (iqs == 0) y[ib].ds = q8_1_ds(d, sum);   // #606: clamped scale/sum (the S26 path bypassed the finite helper)
 }
 template<bool LT, bool LX, int RG, int RD, bool SL = false, bool FQ = false, bool TS = false, int BD = 0>
 void s26_launch_l(const NativeExpertLayout& L, int64_t cap_groups, cudaStream_t s, const unsigned long long* grp_ptr,

src/kernels/cuda/verify_kernels.cu

@@ -4,6 +4,7 @@
 // elementwise.cu) with the same operation order, so a verify window reproduces plain decode bit for bit.
 #include "strata/kernels/verify_kernels.hpp"
 #include "strata/core/emulate.hpp"
+#include "strata/kernels/q8_1_finite.hpp"   // #606: q8_1_ds
 #include "strata/kernels/dp4a.hpp"
 #include "strata/kernels/pdl.hpp"
 
@@ -199,7 +200,7 @@ __device__ __forceinline__ void gdn_q8_1_store(GdnQ81* __restrict__ xq, size_t i
     const float d = amax / 127.0f;
     const int8_t q = amax == 0.0f ? 0 : roundf(xi / d);
     xq[idx / 32].qs[idx % 32] = q;
-    if (idx % 32 == 0) xq[idx / 32].ds = make_half2(d, sum);
+    if (idx % 32 == 0) xq[idx / 32].ds = q8_1_ds(d, sum);   // #606: clamped (the QFUSE path bypassed the finite helper)
 }
 
 template <bool ALL_OUT, bool Q>