window limit: STRATA_WINDOW_MAX_T=12 as a third value (8, 12, 16)

The exact-column instantiations 13..16 of the wide kernels (native MMVQ, the wave variant, the packed Q8_0 and Q6_K launches,
the BF16 multi GEMV) are built only above 12, MAX_NCOLS follows the cap, and kVerifyLongT follows a wide cap (a cap of 12 keeps
--lookup-max-t at 12). Measured for the narrow/wide comparison in phaseA-tests/pr-priorities/RESULTS-NV2.md.

CMakeLists.txt

@@ -99,10 +99,10 @@ option(STRATA_PREFILL_MMQ "Build the optional GGML MMQ prefill path for HIP" OFF
 # the HIP engines (a HIP code object loads on first use, so the extra kernels cost nothing until a wide window forms)
 # and 8 for CUDA, which loads every kernel at start (CUDA_MODULE_LOADING=EAGER: the wide kernels would take VRAM from
 # the expert cache of every card); -DSTRATA_WINDOW_MAX_T=16 builds the wide CUDA engine.
-set(STRATA_WINDOW_MAX_T "auto" CACHE STRING "rows of a verify / batch window: 8, 16 or auto (HIP 16, CUDA 8)")
-set_property(CACHE STRATA_WINDOW_MAX_T PROPERTY STRINGS auto 8 16)
-if(NOT STRATA_WINDOW_MAX_T STREQUAL "auto" AND NOT STRATA_WINDOW_MAX_T STREQUAL "8" AND NOT STRATA_WINDOW_MAX_T STREQUAL "16")
-  message(FATAL_ERROR "STRATA_WINDOW_MAX_T must be auto, 8 or 16 (got '${STRATA_WINDOW_MAX_T}')")
+set(STRATA_WINDOW_MAX_T "auto" CACHE STRING "rows of a verify / batch window: 8, 12, 16 or auto (HIP 16, CUDA 8)")
+set_property(CACHE STRATA_WINDOW_MAX_T PROPERTY STRINGS auto 8 12 16)
+if(NOT STRATA_WINDOW_MAX_T STREQUAL "auto" AND NOT STRATA_WINDOW_MAX_T STREQUAL "8" AND NOT STRATA_WINDOW_MAX_T STREQUAL "12" AND NOT STRATA_WINDOW_MAX_T STREQUAL "16")
+  message(FATAL_ERROR "STRATA_WINDOW_MAX_T must be auto, 8, 12 or 16 (got '${STRATA_WINDOW_MAX_T}')")
 endif()
 if(STRATA_ENABLE_HIP OR STRATA_ENABLE_CUDA)
   if(STRATA_WINDOW_MAX_T STREQUAL "auto")

include/strata/kernels/verify_kernels.hpp

@@ -28,7 +28,7 @@ inline constexpr int kVerifyMaxT = kWindowMaxT;
 /// per-group kernel, the CPU experts and the drafter keep their kVerifyMaxT rows; only the window-wide parts (staging,
 /// head, commit) take more.  A build whose windows already hold 16 rows runs the 16 as one group.
 /// Batch and pipelined windows stay at kVerifyMaxT.
-inline constexpr int kVerifyLongT = kVerifyMaxT >= 16 ? kVerifyMaxT : 2 * kVerifyMaxT;
+inline constexpr int kVerifyLongT = kVerifyMaxT > kVerifyNarrowT ? kVerifyMaxT : 2 * kVerifyMaxT;
 
 /// For token t of T: conv over [history(3) | qkv_0 .. qkv_t] -> SiLU -> L2 norm of the q/k heads -> h[t].
 /// `history` is NOT written, but with `commit` (one token): it then keeps the token, as gdn_conv_commit.  Bitwise

include/strata/kernels/window_limits.hpp

@@ -22,6 +22,6 @@ inline constexpr int kVerifyNarrowT = 8;
 /// The most rows one window can hold in this build (8 or 16).
 inline constexpr int kWindowMaxT = STRATA_VERIFY_MAX_T;
 
-static_assert(kWindowMaxT == 8 || kWindowMaxT == 16, "STRATA_VERIFY_MAX_T is 8 or 16 (the wide kernels are built for 16)");
+static_assert(kWindowMaxT == 8 || kWindowMaxT == 12 || kWindowMaxT == 16, "STRATA_VERIFY_MAX_T is 8, 12 or 16");
 
 }  // namespace strata::kernels

src/core/window_limits_test.cpp

@@ -23,7 +23,7 @@ void check(bool ok, const char* label) {
 int main() {
     using namespace strata::kernels;
     check(kVerifyNarrowT == 8, "the narrow window is the 0.2.7 one: 8 rows");
-    check(kWindowMaxT == 8 || kWindowMaxT == 16, "a window holds 8 or 16 rows");
+    check(kWindowMaxT == 8 || kWindowMaxT == 12 || kWindowMaxT == 16, "a window holds 8, 12 or 16 rows");
     check(kVerifyMaxT == kWindowMaxT, "kVerifyMaxT is the window limit");
     check(kVerifyMaxT >= kVerifyNarrowT, "the wide window is not narrower than the narrow one");
     check(cpu::MAXT == 8, "the CPU kernels' tokens per job stay 8 (a wider window is split into jobs)");

src/kernels/cuda/native_bf16.cu

@@ -382,9 +382,14 @@ void bf16_gemv_fp32_mmvf_multi(const float* x, int64_t ldx, const uint16_t* w, f
     // rows stay with the default case below, the instantiation they always had.
 #if STRATA_VERIFY_MAX_T > 8
 #define STRATA_MMVF_W(N, T) case T: launch_pdl(bf16_f32_mmvf_multi_kernel<N, T, true>, dim3((unsigned) n_out), dim3(N), 0, st, x, ldx, w, y, ldy, (int) n_in, n_tok, mmvf_fast()); break;
+#if STRATA_VERIFY_MAX_T > 12
+#define STRATA_MMVF_W13(N) STRATA_MMVF_W(N, 13) STRATA_MMVF_W(N, 14) STRATA_MMVF_W(N, 15) STRATA_MMVF_W(N, 16)
+#else
+#define STRATA_MMVF_W13(N)
+#endif
 #define STRATA_MMVF_WIDE(N) \
         STRATA_MMVF_W(N, 9) STRATA_MMVF_W(N, 10) STRATA_MMVF_W(N, 11) STRATA_MMVF_W(N, 12) \
-        STRATA_MMVF_W(N, 13) STRATA_MMVF_W(N, 14) STRATA_MMVF_W(N, 15) STRATA_MMVF_W(N, 16)
+        STRATA_MMVF_W13(N)
 #else
 #define STRATA_MMVF_WIDE(N)
 #endif

src/kernels/cuda/native_mmvq.cu

@@ -893,7 +893,7 @@ __global__ void native_small_mmvq_kernel(const Weight* __restrict__ w,
 // compiled only then, so an engine built with STRATA_WINDOW_MAX_T=8 has exactly the kernels it had).  A window has one
 // column per row.  The wide launches read every weight once for all their columns; running 16 columns as two 8-column
 // passes (native_mmvq below does it for more than MAX_NCOLS) reads the dense weights twice and measured slower than 8 rows.
-constexpr int MAX_NCOLS = STRATA_VERIFY_MAX_T > 8 ? 16 : 8;
+constexpr int MAX_NCOLS = STRATA_VERIFY_MAX_T > 8 ? STRATA_VERIFY_MAX_T : 8;
 
 // Each format splits its dot product into `load` (everything that depends only on the weight block: codes,
 // unpacked scales, block scale) and `apply` (the activation loads and the original *_impl expression). `load` runs
@@ -1593,10 +1593,12 @@ void launch_multi(const void* weights, const void* x_q8_1, float* y, int n_in, i
         case 10: launch_multi_n<F, 10>(weights, x_q8_1, y, n_in, n_out, s); break;
         case 11: launch_multi_n<F, 11>(weights, x_q8_1, y, n_in, n_out, s); break;
         case 12: launch_multi_n<F, 12>(weights, x_q8_1, y, n_in, n_out, s); break;
+#if STRATA_VERIFY_MAX_T > 12
         case 13: launch_multi_n<F, 13>(weights, x_q8_1, y, n_in, n_out, s); break;
         case 14: launch_multi_n<F, 14>(weights, x_q8_1, y, n_in, n_out, s); break;
         case 15: launch_multi_n<F, 15>(weights, x_q8_1, y, n_in, n_out, s); break;
         case 16: launch_multi_n<F, 16>(weights, x_q8_1, y, n_in, n_out, s); break;
+#endif
 #endif
         default: throw std::invalid_argument("native MMVQ multi-column launch requires 2 <= ncols <= MAX_NCOLS");
     }
@@ -2406,10 +2408,12 @@ void wave_launch(const void* weights, const void* x_q8_1, float* y, int n_in, in
         case 10: wave_launch_n<F, 10>(weights, x_q8_1, y, n_in, n_out, s); break;
         case 11: wave_launch_n<F, 11>(weights, x_q8_1, y, n_in, n_out, s); break;
         case 12: wave_launch_n<F, 12>(weights, x_q8_1, y, n_in, n_out, s); break;
+#if STRATA_VERIFY_MAX_T > 12
         case 13: wave_launch_n<F, 13>(weights, x_q8_1, y, n_in, n_out, s); break;
         case 14: wave_launch_n<F, 14>(weights, x_q8_1, y, n_in, n_out, s); break;
         case 15: wave_launch_n<F, 15>(weights, x_q8_1, y, n_in, n_out, s); break;
         case 16: wave_launch_n<F, 16>(weights, x_q8_1, y, n_in, n_out, s); break;
+#endif
 #endif
         default: throw std::invalid_argument("native MMVQ (wave) requires 1 <= ncols <= MAX_NCOLS");
     }
@@ -2635,10 +2639,12 @@ bool q8_packed_mmvq(const void* weights, const void* x_q8_1, float* y, int n_in,
         case 10: q8_packed_launch<10>(found->second, x, y, s); break;
         case 11: q8_packed_launch<11>(found->second, x, y, s); break;
         case 12: q8_packed_launch<12>(found->second, x, y, s); break;
+#if STRATA_VERIFY_MAX_T > 12
         case 13: q8_packed_launch<13>(found->second, x, y, s); break;
         case 14: q8_packed_launch<14>(found->second, x, y, s); break;
         case 15: q8_packed_launch<15>(found->second, x, y, s); break;
         case 16: q8_packed_launch<16>(found->second, x, y, s); break;
+#endif
 #endif
         default: throw std::invalid_argument("native MMVQ requires 1 <= ncols <= MAX_NCOLS");
     }
@@ -2780,10 +2786,12 @@ bool q6_packed_mmvq(const void* weights, const void* x_q8_1, float* y, int n_in,
         case 10: q6_packed_launch<10>(found->second, x, y, s); break;
         case 11: q6_packed_launch<11>(found->second, x, y, s); break;
         case 12: q6_packed_launch<12>(found->second, x, y, s); break;
+#if STRATA_VERIFY_MAX_T > 12
         case 13: q6_packed_launch<13>(found->second, x, y, s); break;
         case 14: q6_packed_launch<14>(found->second, x, y, s); break;
         case 15: q6_packed_launch<15>(found->second, x, y, s); break;
         case 16: q6_packed_launch<16>(found->second, x, y, s); break;
+#endif
 #endif
         default: throw std::invalid_argument("native MMVQ requires 1 <= ncols <= MAX_NCOLS");
     }