BW-DENSE: STRATA_SMALL_UNR default 3 and STRATA_GR_V3_PF default 5 on compute capability 12.0 cards (bw_dense_defaults.hpp); STRATA_SM120_PERF=0 / STRATA_DEFAULTS=026 / the variables turn them off

Both are bitwise the loops they replace. Engine A/B: SMALL_UNR on top of the HC change +0.18% agent, +0.28% code_py, 5/5 pairs each, identical text.

include/strata/core/bw_dense_defaults.hpp

@@ -0,0 +1,20 @@
+// strata/core/bw_dense_defaults.hpp - defaults of the bitwise load-ordering variants from the Blackwell Qwen speed work (BW-DENSE).
+//
+// Each is bitwise equal to the kernel it replaces (the screens and the engine A/Bs are in phaseA-tests/blackwell/decisions/E15.md and
+// E21-small-kernel-loads.md), measured on an RTX PRO 5000 Blackwell (compute capability 12.0), so they are on there and nowhere else
+// (unmeasured on other cards).  STRATA_SM120_PERF=0 and STRATA_DEFAULTS=026 turn them off like the rest of the 12.0 set; an explicit variable
+// (STRATA_GR_V3_PF, STRATA_SMALL_UNR) wins on any card in both directions.
+#pragma once
+
+#include "strata/core/defaults_026.hpp"
+#include "strata/core/sm120_series.hpp"
+
+namespace strata {
+
+/// STRATA_GR_V3_PF when unset (gr_down_v3 / gr_up_v3 load-ordering variants): bits 0 (gr_down batched staging loads) and 2 (gr_up batched loads).
+inline int gr_v3_pf_default(int major, int minor) { return sm120_default(major, minor) && !qwen_defaults_026() ? 5 : 0; }
+
+/// STRATA_SMALL_UNR when unset: bit 0 the bf16 router read, bit 1 the GDN a/b projection (constant-trip loops).
+inline int small_unr_default(int major, int minor) { return sm120_default(major, minor) && !qwen_defaults_026() ? 3 : 0; }
+
+}  // namespace strata

src/kernels/cuda/fused_gr.cu

@@ -6,6 +6,7 @@
 #define STRATA_Q81_FASTMATH_DIV 1   // #3948: the q8_1 tails divide as native_mmvq.cu (--use_fast_math) does, see q8_1_finite.hpp
 #include "strata/kernels/q8_1_finite.hpp"   // #606: q8_1_ds
 #include "strata/core/geometry.hpp"
+#include "strata/core/bw_dense_defaults.hpp"
 #include "strata/kernels/fused_gr.hpp"
 #include "strata/kernels/bf16_bits.hpp"
 #include "strata/kernels/verify_kernels.hpp"
@@ -2551,7 +2552,7 @@ static int v3_pf_mask() {
         const bool ok = cudaDeviceGetAttribute(&ma, cudaDevAttrComputeCapabilityMajor, dev) == cudaSuccess &&
                         cudaDeviceGetAttribute(&mi, cudaDevAttrComputeCapabilityMinor, dev) == cudaSuccess;
         if (!ok) cudaGetLastError();
-        per_dev[dev] = 1 + (ok && ma == 12 && mi == 0 ? 5 : 0);
+        per_dev[dev] = 1 + (ok ? strata::gr_v3_pf_default(ma, mi) : 0);
     }
     return per_dev[dev] - 1;
 #endif

src/kernels/cuda/native_bf16.cu

@@ -1,3 +1,4 @@
+#include "strata/core/bw_dense_defaults.hpp"
 #include "strata/kernels/bf16_gemv.hpp"
 #include "strata/kernels/window_limits.hpp"
 #include "s26_tsum.cuh"
@@ -278,10 +279,25 @@ void launch_rows(const float* x, int64_t ldx, const uint16_t* w, float* y, int64
     else bf16_f32_mmvf_rows_kernel<B, 8, RPB><<<nb, B, 0, st>>>(x, ldx, w, y, ldy, n_in, n_out, n_tok);
 }
 
-// BW-DENSE STRATA_SMALL_UNR bit 0: see bf16_f32_mmvf_multi_kernel (default 0 until measured)
+// BW-DENSE STRATA_SMALL_UNR bit 0: see bf16_f32_mmvf_multi_kernel (bw_dense_defaults.hpp: on for a 12.0 card, else off; the variable wins)
 int mmvf_fast() {
-    static const int v = [] { const char* e = std::getenv("STRATA_SMALL_UNR"); return e != nullptr ? (std::atoi(e) & 1) : 0; }();
-    return v;
+    static int per_dev[64];   // 0 unknown, else bit + 1
+    const char* e = std::getenv("STRATA_SMALL_UNR");
+    if (e != nullptr) return std::atoi(e) & 1;
+#if defined(__HIPCC__)
+    return 0;
+#else
+    int dev = 0;
+    if (cudaGetDevice(&dev) != cudaSuccess || dev < 0 || dev >= 64) { cudaGetLastError(); return 0; }
+    if (!per_dev[dev]) {
+        int ma = 0, mi = 0;
+        const bool ok = cudaDeviceGetAttribute(&ma, cudaDevAttrComputeCapabilityMajor, dev) == cudaSuccess &&
+                        cudaDeviceGetAttribute(&mi, cudaDevAttrComputeCapabilityMinor, dev) == cudaSuccess;
+        if (!ok) cudaGetLastError();
+        per_dev[dev] = 1 + (ok ? (strata::small_unr_default(ma, mi) & 1) : 0);
+    }
+    return per_dev[dev] - 1;
+#endif
 }
 int mmvf_block_size(int64_t n_in) {
     int best = 32;

src/kernels/cuda/verify_kernels.cu

@@ -3,6 +3,7 @@
 // The per-token arithmetic of every kernel here is transcribed from its single-token original (fused_gdn.cu,
 // 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/bw_dense_defaults.hpp"
 #include "strata/core/geometry.hpp"
 #include "strata/core/emulate.hpp"
 #define STRATA_Q81_FASTMATH_DIV 1   // #3948: the q8_1 stores divide as native_mmvq.cu (--use_fast_math) does, see q8_1_finite.hpp
@@ -1141,10 +1142,25 @@ bool gdn_commit_layers(float* state, int64_t state_stride, int64_t conv_off, con
     return true;
 }
 
-// BW-DENSE STRATA_SMALL_UNR bit 1: see gdn_ab_body (default 0 until measured)
+// BW-DENSE STRATA_SMALL_UNR bit 1: see gdn_ab_body (bw_dense_defaults.hpp: on for a 12.0 card, else off; the variable wins)
 static int gdn_ab_fast() {
-    static const int v = [] { const char* e = std::getenv("STRATA_SMALL_UNR"); return e != nullptr ? ((std::atoi(e) >> 1) & 1) : 0; }();
-    return v;
+    static int per_dev[64];   // 0 unknown, else bit + 1
+    const char* e = std::getenv("STRATA_SMALL_UNR");
+    if (e != nullptr) return (std::atoi(e) >> 1) & 1;
+#if defined(__HIPCC__)
+    return 0;
+#else
+    int dev = 0;
+    if (cudaGetDevice(&dev) != cudaSuccess || dev < 0 || dev >= 64) { cudaGetLastError(); return 0; }
+    if (!per_dev[dev]) {
+        int ma = 0, mi = 0;
+        const bool ok = cudaDeviceGetAttribute(&ma, cudaDevAttrComputeCapabilityMajor, dev) == cudaSuccess &&
+                        cudaDeviceGetAttribute(&mi, cudaDevAttrComputeCapabilityMinor, dev) == cudaSuccess;
+        if (!ok) cudaGetLastError();
+        per_dev[dev] = 1 + (ok ? ((strata::small_unr_default(ma, mi) >> 1) & 1) : 0);
+    }
+    return per_dev[dev] - 1;
+#endif
 }
 void gdn_ab_multi(const float* x, const uint16_t* w_alpha, const uint16_t* w_beta, const float* dt, const float* ssm_a,
                   float* gate, float* beta, int n_embd, int h_v, int n_tok, void* stream) {