tests: gdn_rec_head also compares the PP (variant 4) and two-column (variant 5) recurrence kernels, bitwise

T = 1, 7, 8, 9, 17, 100, 2048, 2051 (full and partial chunks), y, its FP16 bits and the final state: 0 values differ.

src/prefill/kernels.cu

@@ -1374,6 +1374,18 @@ void gdn_recurrence_variant(int variant, float* state, const float* h, const flo
         }
         return;
     }
+    if (variant == 4 || variant == 5) {   // diagnostics: the quad recurrence's PP / two-column kernels + the norm
+        if (T > 0) {
+            if (variant == 5)
+                gdn_rec_quad2c_kernel<<<HV * (S / C2CB), QTH, 0, (cudaStream_t) stream>>>(state, h, gate, beta, y, T);
+            else
+                gdn_rec_quad_pp_kernel<<<HV * (S / QCB), QTH, 0, (cudaStream_t) stream>>>(state, h, gate, beta, y, T);
+            gdn_out_norm_loop_kernel<true><<<(unsigned) std::min<int64_t>(T * HV, 4096), S, 0, (cudaStream_t) stream>>>(
+                z, gamma, eps, y, y16, T * HV, ld16);
+        }
+        check("gdn_recurrence (quad pp)");
+        return;
+    }
     if (variant == 1) {
         if (T > 0) {
             static const int pp = [] { const char* v = std::getenv("STRATA_GDN_PP"); return v ? std::atoi(v) : 0; }();

tests/hip/gdn_rec_head.cpp

@@ -61,7 +61,7 @@ int main() {
         const float* dgm = dev(gamma);
         std::vector<float> y_ref, st_ref;
         std::vector<uint16_t> y16_ref;
-        for (int variant : {0, 3, 1}) {
+        for (int variant : {0, 3, 1, 4, 5}) {
             float* dst = dev(state);
             float* dy = nullptr;
             uint16_t* dy16 = nullptr;