perf(GFX1100-TG200): T4a MMVQ K-quant decode GEMV arm

CMakeLists.txt

@@ -1738,7 +1738,8 @@ if(VLLM_CPP_HIP)
     src/vt/rocm/rocm_mla_ops.hip
     src/vt/rocm/rocm_mla_attn.hip
     src/vt/rocm/rocm_skinny_gemm.hip
-    src/vt/rocm/rocm_ops.hip)
+    src/vt/rocm/rocm_ops.hip
+    src/vt/rocm/rocm_quant_dot.hip)
   if(VLLM_CPP_HIP_ARCHITECTURES)
     set_source_files_properties(
       src/vt/rocm/rocm_backend.hip
@@ -1767,6 +1768,7 @@ if(VLLM_CPP_HIP)
       src/vt/rocm/rocm_mla_attn.hip
       src/vt/rocm/rocm_skinny_gemm.hip
       src/vt/rocm/rocm_ops.hip
+      src/vt/rocm/rocm_quant_dot.hip
       PROPERTIES HIP_ARCHITECTURES "${VLLM_CPP_HIP_ARCHITECTURES}")
   endif()
   # Prefer the absolute path inside ${ROCM_PATH}/lib, fall back to the bare name,

docs/ENVIRONMENT.md

@@ -224,6 +224,8 @@ portable/reference path. In normal operation leave them unset.
 | `VT_GPU_SAMPLE` | on (CUDA) | Host-side sampling instead of on-GPU sampling |
 | `VT_ROCM_Q8K_BLOCK` | unset (cooperative only on queue-device-resolved `gfx1100`) | Exact `0` selects the permanent legacy A/B arm. Exact `1` forces the cooperative candidate diagnostically. Explicit `1` outside validated `gfx1100` is not a support claim or a default claim. Unset keeps the legacy arm on `gfx1200`, `gfx1201`, unknown architectures, and architecture-resolution failure. Every other value is refused. See [Q8_K activation quantization](ROCM.md#select-q8_k-activation-quantization) |
 | `VT_ROCM_QUANT_WMMA` | on | ROCm keep-quant Q6_K and Q4_K prefill GEMM: the RDNA4 rocWMMA int8 tile arm (`KQuantGemmKWmmaQ6K`, `KQuantGemmKWmmaQ4K`), queue-device-resolved `gfx1200`/`gfx1201` only and only when `m >= 16 && n >= 16` (at least one full 16-wide tile in each dimension, not exact alignment — the scalar `KQuantGemmK` remaps its launched index to enumerate only the remainder the WMMA corner leaves untouched (the bottom strip past the aligned row boundary, then the right strip past the aligned column boundary), so a non-16-multiple M/N still takes the WMMA arm over its floor(M/16)xfloor(N/16) corner without paying a full-`m*n`-grid launch for the fill). `0` forces the scalar `Dp4a` arm this row's spec (`KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA4`, issue #2109) is chasing a performance gap against, on every architecture and every shape. Bit-identical to the scalar arm by construction (the WMMA tile's raw int8 dot is scaled and reduced in the same integer arithmetic the scalar path uses; only the one f32 scale product per superblock is shared with it) — this is a same-binary A/B, not a correctness fallback |
+| `VT_GEMV_MMVQ` | off | `=1` routes the K-quant (Q4_K/Q5_K/Q6_K) `MatmulBTQuant` decode (m==1 only) through the MMVQ-style GEMV arm in `rocm_grouped_gemm.hip`, which folds activation quantization into the GEMV prologue; default keeps the baseline `KQuantGemmK` path byte-unchanged. The arm is bit-equal to the baseline output at every (Fmt, nsb, j); the fold crossover is tunable with `VT_GEMV_MMVQ_FOLD_MAX`. Flag read per call, so a captured decode graph picks the arm up at capture time |
+| `VT_SKINNY_BF16` | off | `=1` dispatches the f32-output decode-skinny arm (`WvSplitKBTToF32`) for the bf16-input GEMVs whose consumers want f32 — the Qwen3.5 GDN BA projections' g/beta chain — instead of rocBLAS's large-M tile |
 | `VT_GDN_PACKED_DECODE` | on (CUDA GDN) | Unpacked GDN decode path |
 | `VT_GDN_DECODE_BV` | `32` (CUDA GDN decode experiment) | Exact `16` selects the byte-identical 16-value fused-recurrence tile; unset and every other spelling keep the 32-value schedule. Experimental opt-in; no release or cross-hardware default change |
 | `VT_GDN_DECODE_SWIZZLE` | `0` (CUDA GDN decode experiment) | Exact `1` enables the shared-memory bank swizzle only for the `BV=16`, `Dv=Dk=128`, eight-lane production geometry; all other values and shapes keep the incumbent layout |

src/vt/rocm/rocm_matmul_hipblaslt.hip

@@ -17,6 +17,7 @@
 #include <cstdint>
 #include <cstdio>
 #include <cstdlib>
+#include <atomic>
 #include <cstring>
 #include <mutex>
 #include <stdexcept>
@@ -529,11 +530,49 @@ void MatmulKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tensor& b) {
             "hipblasGemmEx NN");
 }
 
+// Host-side routing witness for the bf16-in/f32-out decode population (the
+// Qwen3.5 GDN BA pair; evidence 15.1): process-global counters bumped on
+// exactly the branch taken per dispatch. Both routes are numerically valid,
+// so outputs alone cannot witness routing. Same shape as the T4a MMVQ
+// counters (rocm_grouped_gemm.hip).
+struct SkinnyF32RouteCounts {
+  long long blas;    // fell through to hipblasGemmEx (default route)
+  long long skinny;  // took the VT_SKINNY_BF16 wvSplitK-class arm
+};
+
+namespace {
+std::atomic<long long> g_skinny_f32_route_blas{0};
+std::atomic<long long> g_skinny_f32_route_skinny{0};
+}  // namespace
+
+void SkinnyF32ResetRouteCountsForTesting() {
+  g_skinny_f32_route_blas.store(0, std::memory_order_relaxed);
+  g_skinny_f32_route_skinny.store(0, std::memory_order_relaxed);
+}
+
+SkinnyF32RouteCounts SkinnyF32RouteCountsForTesting() {
+  return {g_skinny_f32_route_blas.load(std::memory_order_relaxed),
+          g_skinny_f32_route_skinny.load(std::memory_order_relaxed)};
+}
+
+// Lever B2 opt-in arm (evidence 15.1): VT_SKINNY_BF16=1 serves bf16-in/
+// f32-out decode-skinny shapes with the wvSplitK geometry instead of
+// rocBLAS's large-M tile. Read PER CALL (cuda_quant_dot.cu convention) so
+// in-process tests and captured graphs pick the arm up at launch time.
+// Default OFF: the default path is byte-unchanged.
+bool SkinnyBf16F32OutEnabled() {
+  if (const char* e = std::getenv("VT_SKINNY_BF16")) return e[0] == '1';
+  return false;
+}
+
 // out[M,N] = a[M,K] @ b[N,K]^T
 // wvSplitK skinny-GEMM host entry (rocm_skinny_gemm.hip, #487). External
 // vt::rocm linkage to match the definition; declared beside its only caller.
 void WvSplitKBT(hipStream_t s, void* out, const void* a, const void* b, int M, int N,
                 int K, int device);
+// Lever B2 f32-output variant of the same kernel geometry.
+void WvSplitKBTToF32(hipStream_t s, void* out, const void* a, const void* b,
+                     int M, int N, int K, int device);
 
 // Row-major trick: gemm(OP_T, OP_N, N, M, K, B, K, A, a_rs, C, N)
 //   BLAS: C = op(A)*op(B) with opA=T => A is KxN in col form = row B[N,K]
@@ -583,6 +622,25 @@ void MatmulBTKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tensor& b)
     return;
   }
 
+  // Lever B2 (evidence 15.1): bf16-in/f32-out decode-skinny — the Qwen3.5
+  // GDN BA projections (N=32, K=2560, m=1) emit f32 and today starve on
+  // rocBLAS's large-M MT128x32x16 tile (~73.7us for a 164 KiB weight).
+  // Same donor guards as the bf16 arm above (the kernel is the same
+  // geometry, so the same tail-safety constraints apply), opt-in via
+  // VT_SKINNY_BF16=1, default OFF and byte-unchanged.
+  const bool skinny_f32_pop =
+      bf16 && out.dtype == DType::kF32 && M >= 1 && M <= 4;
+  if (skinny_f32_pop && SkinnyBf16F32OutEnabled() && (K % 8) == 0 && N > 8 &&
+      (N % 2) == 0 && a.stride[0] == K && K * M <= 32768 &&
+      vt::rocm::SkinnyGemmArchOk(q.device.index, vt::rocm::DeviceArchName)) {
+    WvSplitKBTToF32(s, out.data, a.data, b.data, static_cast<int>(M),
+                    static_cast<int>(N), static_cast<int>(K), q.device.index);
+    g_skinny_f32_route_skinny.fetch_add(1, std::memory_order_relaxed);
+    return;
+  }
+  if (skinny_f32_pop)
+    g_skinny_f32_route_blas.fetch_add(1, std::memory_order_relaxed);
+
   // Decode: M=1 BF16 GEMV
   if (M == 1 && bf16 && out.dtype == DType::kBF16 && a.stride[0] == K && GemvEnabled()) {
     Bf16GemvBT(s, out.data, a.data, b.data, static_cast<int>(N), static_cast<int>(K), 1.f, 0.f);