#606: the fused SwiGLU q8_1 quantizers keep the block's scale and sum finite too

0.1.39 clamped two q8_1 activation quantizers (#606: an fp16 sum or scale past 65504 rounds to inf, the dot products
then read inf * 0 = NaN, and a NaN in the residual ends as one token repeated). Two more store the block unclamped:
- native_swiglu_quantize_q8_1_kernel (native_mmvq.cu), the shared expert's fused SwiGLU + q8_1, on by default
  (STRATA_FUSED_SWIGLU_Q81) in decode and the multi-token pass;
- native_gu_fused_kernel (iq_kernels.cu), the routed experts' fused SwiGLU + q8_1 (gfx906 builds, STRATA_EXP_MODE=8).
Both now go through q8_1_finite / q8_1_quant / q8_1_ds as their siblings do: every finite block is stored bit for bit
as before. iq_multi_parity's #606 check covers the shared expert's one (gate = 32, up = x / 32: its activations are x
exactly). rwkeyes found the first one (#606).

src/kernels/cuda/iq_kernels.cu

@@ -2413,11 +2413,11 @@ __global__ void __launch_bounds__(256) native_gu_fused_kernel(const unsigned lon
                 amax = fmaxf(amax, __shfl_xor_sync(0xffffffffu, amax, o));
                 sum += __shfl_xor_sync(0xffffffffu, sum, o);
             }
-            const float d = amax / 127.0f;
-            const int8_t q = amax == 0.0f ? 0 : roundf(xi / d);
+            const float d = q8_1_finite(amax / 127.0f);   // #606, as q8_1_store: finite blocks bit for bit
+            const int8_t q = q8_1_quant(xi, d, amax);
             block_q8_1* y = hq + (size_t) (c0 + warp) * hb + blockIdx.x;
             y->qs[lane] = q;
-            if (lane == 0) y->ds = make_half2(d, sum);
+            if (lane == 0) y->ds = q8_1_ds(d, sum);
         }
     }
 }

src/kernels/cuda/native_mmvq.cu

@@ -163,10 +163,10 @@ __global__ void native_swiglu_quantize_q8_1_kernel(const float* __restrict__ gat
     const float xi = __fmul_rn(__fdividef(gi, __fadd_rn(1.0f, __expf(-gi))), up[i]);
     const float amax = warp_max(fabsf(xi));
     const float sum = warp_sum(xi);
-    const float d = amax / 127.0f;
-    const int8_t q = amax == 0.0f ? 0 : roundf(xi / d);
+    const float d = q8_1_finite(amax / 127.0f);   // #606, as native_quantize_q8_1_kernel: finite blocks bit for bit
+    const int8_t q = q8_1_quant(xi, d, amax);
     y[i / Q8K].qs[i % Q8K] = q;
-    if (i % Q8K == 0) y[i / Q8K].ds = make_half2(d, sum);
+    if (i % Q8K == 0) y[i / Q8K].ds = q8_1_ds(d, sum);
 }
 
 // Exact pinned vec_dot_q5_K_q8_1_impl_vmmq expression and integer dot order.

src/kernels/iq_multi_parity.cpp

@@ -326,9 +326,10 @@ void bench(cudaStream_t s, std::mt19937& rng) {
     }
 }
 
-// #606: both q8_1 activation quantizers (quantize_q8_1_rows and native_quantize_q8_1) against a host transcription
-// of the 0.1.38 formula: every finite block bit for bit; a block whose sum (or scale) overflowed fp16 is now stored
-// finite (the largest half, its sign), with its values in the int8 range.
+// #606: the q8_1 activation quantizers (quantize_q8_1_rows, native_quantize_q8_1 and the shared expert's fused
+// native_swiglu_quantize_q8_1) against a host transcription of the 0.1.38 formula: every finite block bit for bit; a
+// block whose sum (or scale) overflowed fp16 is now stored finite (the largest half, its sign), with its values in the
+// int8 range.  The fused one gets gate = 32 and up = x / 32: silu(32) is 32 in float, so its activations are x exactly.
 void check_q8_1_finite(cudaStream_t s, std::mt19937& rng) {
     const int n = 2560, rows = 3, nb = n / 32;
     std::vector<float> x = random_x((size_t) rows * n, rng);
@@ -338,10 +339,17 @@ void check_q8_1_finite(cudaStream_t s, std::mt19937& rng) {
     x[(size_t) n + 40] = 2.0e7f;                                                    // amax / 127 overflowed too
     float* dx = dalloc<float>(x.size());
     ck(cudaMemcpy(dx, x.data(), x.size() * 4, cudaMemcpyHostToDevice), "x");
+    std::vector<float> gate(x.size(), 32.0f), up(x.size());
+    for (size_t i = 0; i < x.size(); ++i) up[i] = x[i] / 32.0f;
+    float* dg = dalloc<float>(x.size());
+    float* du = dalloc<float>(x.size());
+    ck(cudaMemcpy(dg, gate.data(), x.size() * 4, cudaMemcpyHostToDevice), "gate");
+    ck(cudaMemcpy(du, up.data(), x.size() * 4, cudaMemcpyHostToDevice), "up");
     uint8_t* dq = dalloc<uint8_t>((size_t) rows * nb * 36);
-    for (int path = 0; path < 2; ++path) {
+    for (int path = 0; path < 3; ++path) {
         if (path == 0) k::quantize_q8_1_rows(dx, rows, n, dq, s);
-        else k::native_quantize_q8_1(dx, dq, n, rows, s);
+        else if (path == 1) k::native_quantize_q8_1(dx, dq, n, rows, s);
+        else k::native_swiglu_quantize_q8_1(dg, du, dq, n, rows, s);
         ck(cudaStreamSynchronize(s), "quantize");
         std::vector<uint8_t> got((size_t) rows * nb * 36);
         ck(cudaMemcpy(got.data(), dq, got.size(), cudaMemcpyDeviceToHost), "q");
@@ -370,7 +378,7 @@ void check_q8_1_finite(cudaStream_t s, std::mt19937& rng) {
                     want[4 + i] = (uint8_t) (int8_t) (amax == 0.0f ? 0.0f : std::round(v[i] / d));
                 // native_mmvq.cu is built with --use_fast_math (its divisions are approximate): there the sum must be
                 // the same bits, the scale within one fp16 step and each value within one step
-                bool near = path == 1 && gs == os && std::abs((int) gd - (int) od) <= 1;
+                bool near = path >= 1 && gs == os && std::abs((int) gd - (int) od) <= 1;
                 for (int i = 0; near && i < 32; ++i) near = std::abs((int) (int8_t) g[4 + i] - (int) (int8_t) want[4 + i]) <= 1;
                 if (std::memcmp(want, g, 36) == 0 || near) ++same;
                 else { ++bad; std::printf("  q8_1 path %d block %d: changed although it was finite\n", path, b); }
@@ -384,10 +392,13 @@ void check_q8_1_finite(cudaStream_t s, std::mt19937& rng) {
             }
         }
         std::printf("q8_1 finite (%s): %d finite blocks as before, %d overflowed blocks clamped finite%s\n",
-                    path == 0 ? "quantize_q8_1_rows" : "native_quantize_q8_1", same, clamped, bad ? "  FAIL" : "");
+                    path == 0 ? "quantize_q8_1_rows" : path == 1 ? "native_quantize_q8_1" : "native_swiglu_quantize_q8_1",
+                    same, clamped, bad ? "  FAIL" : "");
         g_fail += bad + (clamped != 3);
     }
     cudaFree(dx);
+    cudaFree(dg);
+    cudaFree(du);
     cudaFree(dq);
 }