feat(GFX1100-TG200): T21 keep-quant for V-head row-permuted GDN projections

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,13 @@ 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_ROWPERM_KEEP_QUANT` | off | `=1` keeps the row-permuted V-head GDN projections (in the tiled order the row permutation produces) as K-quant instead of expanding to bf16 at load; the runtime gather supplies the permutation. Opt-in; the default reorders then expands |
+| `VT_RMSNORM_ROW_COOP` | off | `=1` selects the cooperative single-row RMSNorm arm (two wavefront `shfl_down` trees plus one cross-wavefront shared-memory combine — two `__syncthreads` instead of nine) when h is a multiple of 256 and at most one block wide; the float reduction order CHANGES, so this is an opt-in riding the campaign config, not a default change. Eligibility falls back to the ported body automatically |
+| `VT_RMSNORM_LDS_QUANT` | on (when `VT_RMSNORM_ROW_COOP=1` and a Q8_K output is active) | `=0` disables the T24 LDS-buffered quant epilogue that stores the normalized row to shared memory during pass 3 so the Q8_K quant epilogue reads from LDS instead of re-reading global memory; the quantized bytes are unchanged either way |
+| `VT_QUANT_Q8K_WARP` | off | `=1` selects the warp-cooperative `QuantizeQ8KKWarpCoop` activation quantizer (16 super-blocks per 128-thread block, per-chunk ascending amax with cross-warp shuffle reduction) instead of the selector-dispatched default on the decode path; byte-identical output, asserted on tied-amax rows |
+| `VT_NORM_QUANT_FUSED` | off | `=1` fuses the Q8_K activation-superblock quantizer into the cooperative RmsNorm epilogue (one launch instead of norm + standalone quant); the producer registers its scratch in a consumer token bridge, so the consuming `MatmulBTQuant` decode SKIPS the standalone quant launch when the row set matches. Byte-equal to the standalone path by construction (both instantiate the shared `QuantQ8KSBlock` body from `rocm_act_quant.h`) and asserted op-level |
 | `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/vllm/model_executor/models/qwen3_5_gguf_weights.cpp

@@ -382,6 +382,23 @@ void ReorderVRows(std::vector<T>& buf, int64_t cols, int64_t row_off,
   }
   std::memcpy(base, seg.data(), seg.size() * sizeof(T));
 }
+// Pointer-based overload for OwnedBytes (T=uint8_t, sizeof(T)=1).
+void ReorderVRows(uint8_t* buf, int64_t cols, int64_t row_off,
+                  int64_t num_k, int64_t num_v_per_k, int64_t head_rows) {
+  const int64_t num_v = num_k * num_v_per_k;
+  const int64_t head_stride = head_rows * cols;
+  std::vector<uint8_t> seg(static_cast<size_t>(num_v) * head_stride);
+  uint8_t* base = buf + row_off * cols;
+  for (int64_t k = 0; k < num_k; ++k) {
+    for (int64_t r = 0; r < num_v_per_k; ++r) {
+      const int64_t g = k * num_v_per_k + r;
+      const int64_t t = r * num_k + k;
+      std::memcpy(seg.data() + g * head_stride, base + t * head_stride,
+                  static_cast<size_t>(head_stride));
+    }
+  }
+  std::memcpy(base, seg.data(), seg.size());
+}
 
 // Reorder the full column range [0, cols) of a [rows, cols] row-major buffer
 // (cols = num_v * head_cols) from GGUF tiled to HF grouped order (out_proj).
@@ -1072,14 +1089,30 @@ GdnLayerWeights LoadGdnGguf(const GgufFile& g, int64_t il, const HfConfig& c,
   const int64_t key_dim = num_k * c.linear_key_head_dim;
   const bool reorder = num_v != num_k && num_k > 0 && (num_v % num_k) == 0;
   const int64_t rpk = num_k > 0 ? num_v / num_k : 1;  // num_v_per_k
-  // When the V-head reorder is active these projections are LAYOUT-rewritten
-  // at load, so they are kTransformedWeight and can never keep their blocks;
-  // without it they are ordinary verbatim GEMM weights. (out_proj's reorder
-  // permutes COLUMNS, which live inside a block, so it is unconditionally
-  // block-unsafe when active — same rule, stated per tensor below.)
+  // When the V-head reorder is active, the projections are LAYOUT-rewritten at
+  // load. For COLUMN-permuted tensors (out_proj/ssm_out) the reorder cuts across
+  // quantization block boundaries, so they are kTransformedWeight and must
+  // expand to bf16. For ROW-permuted tensors (in_proj_qkv, in_proj_z) the
+  // reorder only changes row order — quantization blocks are along the K
+  // (column) dimension and are self-contained per row — so the blocks can be
+  // kept and the permutation applied to the block rows at load time (T21).
+  // Without reorder they are ordinary verbatim GEMM weights. Column-permuted
+  // tensors (out_proj/ssm_out) stay kTransformedWeight and expand to bf16.
+  // T21 env gate: VT_GDN_ROWPERM_KEEP_QUANT=0 forces the row-permuted tensors
+  // back to kTransformedWeight (bf16 expansion) for A/B isolation.
+  // Opt-in, by VALUE like its VT_GDN_COLPERM_KEEP_QUANT sibling: unset keeps
+  // the load path byte-identical to the trunk (reorder, then expand to bf16).
+  // A default-on flip here silently changed the trunk numerics of every
+  // row-permuted V-head projection and broke the token-exact gate -- caught on
+  // the gfx1100 engine leg, reverted to opt-in.
+  const char* rpkq = std::getenv("VT_GDN_ROWPERM_KEEP_QUANT");
+  const bool rowperm_keep =
+      rpkq != nullptr && rpkq[0] == '1' && rpkq[1] == '\0';
   const GgufTensorRole proj_role = reorder
                                        ? GgufTensorRole::kTransformedWeight
                                        : GgufTensorRole::kMatmulWeight;
+  const GgufTensorRole rowperm_role =
+      (reorder && rowperm_keep) ? GgufTensorRole::kMatmulWeight : proj_role;
   // GdnLayerWeights carries an Nvfp4Weight ONLY for out_proj, and even that is
   // unreachable on the 27B because the V-column reorder makes ssm_out
   // kTransformedWeight. The in_proj family has no fp4 field at all. So the GDN
@@ -1092,11 +1125,27 @@ GdnLayerWeights LoadGdnGguf(const GgufFile& g, int64_t il, const HfConfig& c,
   GdnLayerWeights gdn;
 
   // in_proj_qkv <- attn_qkv [conv_dim, H]; only the trailing V rows reorder.
+  // T21: ReorderVRows is a row permutation (block-safe for K-quant). Route as
+  // kMatmulWeight to allow keep-quant, then permute the block rows in place.
+  // Saves ~661 MB/tok of bf16 read amplification (24 Q5_K tensors × 2.9x).
+  // The forward pass already dispatches quantized nk=true weights through
+  // vt::MatmulBT → matmul_bt_quant, so no forward-pass change is needed.
   {
     const std::string nm = Blk(il, "attn_qkv.weight");
-    const GgufResidency r = pol.Route(g.Get(nm), proj_role);
-    if (r != GgufResidency::kExpandBf16) {
-      const GgufTensorInfo& ti = g.Get(nm);
+    const GgufTensorInfo& ti = g.Get(nm);
+    const GgufResidency r = pol.Route(ti, rowperm_role);
+    if (r == GgufResidency::kKeepQuant) {
+      // Force a copy (not mmap) so the block rows can be permuted in place.
+      OwnedTensor qk = OwnGgufQuantBlocks(ti, ti.shape[0], ti.shape[1], 0,
+                                          /*mmap_src=*/nullptr);
+      if (reorder) {
+        const int64_t row_bytes = static_cast<int64_t>(qk.bytes.size()) /
+                                  ti.shape[0];
+        ReorderVRows(qk.bytes.data(), row_bytes, /*row_off=*/2 * key_dim,
+                     num_k, rpk, dv);
+      }
+      gdn.in_proj_qkv = std::move(qk);
+    } else if (r != GgufResidency::kExpandBf16) {
       gdn.in_proj_qkv =
           OwnGgufKeptSlice(g, pol, ti, r, ti.shape[0], ti.shape[1], 0);
     } else {
@@ -1109,11 +1158,22 @@ GdnLayerWeights LoadGdnGguf(const GgufFile& g, int64_t il, const HfConfig& c,
     }
   }
   // in_proj_z <- attn_gate [value_dim, H]; all rows are V.
+  // T21: Same row-permutation keep-quant path as in_proj_qkv above.
+  // Saves ~360 MB/tok of bf16 read amplification (24 Q4_K tensors × 2.9x).
   {
     const std::string nm = Blk(il, "attn_gate.weight");
-    const GgufResidency r = pol.Route(g.Get(nm), proj_role);
-    if (r != GgufResidency::kExpandBf16) {
-      const GgufTensorInfo& ti = g.Get(nm);
+    const GgufTensorInfo& ti = g.Get(nm);
+    const GgufResidency r = pol.Route(ti, rowperm_role);
+    if (r == GgufResidency::kKeepQuant) {
+      OwnedTensor qk = OwnGgufQuantBlocks(ti, ti.shape[0], ti.shape[1], 0,
+                                          /*mmap_src=*/nullptr);
+      if (reorder) {
+        const int64_t row_bytes = static_cast<int64_t>(qk.bytes.size()) /
+                                  ti.shape[0];
+        ReorderVRows(qk.bytes.data(), row_bytes, 0, num_k, rpk, dv);
+      }
+      gdn.in_proj_z = std::move(qk);
+    } else if (r != GgufResidency::kExpandBf16) {
       gdn.in_proj_z =
           OwnGgufKeptSlice(g, pol, ti, r, ti.shape[0], ti.shape[1], 0);
     } else {

src/vt/rocm/rocm_act_quant.h

@@ -0,0 +1,124 @@
+// Shared ROCm device-side Q8_K activation-superblock quantizer (Lever C,
+// GFX1100-TG200-NORMQ). One source of truth for the byte-exactness-critical
+// numeric path: src/vt/rocm/rocm_grouped_gemm.hip (standalone QuantizeQ8KK +
+// MMVQ fused-prologue) AND src/vt/rocm/rocm_rmsnorm.hip (producer-fused
+// epilogue behind VT_NORM_QUANT_FUSED=1) both instantiate THIS body, so
+// "byte-equal vs standalone" holds by construction rather than by two copies
+// drifting. Contract carried over from cuda_quant_dot.cu QuantizeQ8KPreqKernel:
+// the amax carries its ORIGINAL element index and ties break by LOWEST index
+// (`ax > amax`, never `>=`); tests assert this on tied-amax rows.
+//
+// The helpers here were moved verbatim out of rocm_grouped_gemm.hip's
+// anonymous namespace (clean cutover, no second copy left behind); every
+// consumer in that file keeps resolving the same names through this include.
+#ifndef VLLM_CPP_SRC_VT_ROCM_ROCM_ACT_QUANT_H_
+#define VLLM_CPP_SRC_VT_ROCM_ROCM_ACT_QUANT_H_
+
+#include <cstdint>
+#include "vt/dtype.h"
+#include "vt/cpu/cpu_quant_blocks.h"
+
+namespace vt::rocm {
+
+enum class ActDT : int { kF32 = 0, kF16 = 1, kBF16 = 2 };
+
+inline ActDT ActDtOf(DType dt) {
+  return dt == DType::kF32 ? ActDT::kF32 : dt == DType::kF16 ? ActDT::kF16 : ActDT::kBF16;
+}
+
+__device__ inline float DF16ToF32(uint16_t h) {
+  uint32_t sign = static_cast<uint32_t>(h & 0x8000) << 16;
+  uint32_t exp = (h >> 10) & 0x1F;
+  uint32_t mant = h & 0x3FF;
+  if (exp == 0x1F) return __int_as_float(sign | 0x7F800000 | (mant << 13));
+  if (exp == 0) {
+    if (mant == 0) return __int_as_float(sign);
+    int shift = 0;
+    while ((mant & 0x400) == 0) { mant <<= 1; ++shift; }
+    mant &= 0x3FF;
+    return __int_as_float(sign | ((113 - shift) << 23) | (mant << 13));
+  }
+  return __int_as_float(sign | ((exp + 112) << 23) | (mant << 13));
+}
+__device__ inline float DBF16ToF32(uint16_t b) {
+  return __int_as_float(static_cast<uint32_t>(b) << 16);
+}
+__device__ inline uint16_t DF32ToBF16(float f) {
+  uint32_t u = __float_as_int(f);
+  if ((u & 0x7F800000) == 0x7F800000 && (u & 0x7FFFFF))
+    return static_cast<uint16_t>((u >> 16) | 0x0040);
+  uint32_t rounding = 0x7FFF + ((u >> 16) & 1);
+  return static_cast<uint16_t>((u + rounding) >> 16);
+}
+__device__ inline uint16_t DF32ToF16(float f) {
+  uint32_t u = __float_as_uint(f);
+  uint16_t sign = static_cast<uint16_t>((u >> 16) & 0x8000);
+  int32_t exp = static_cast<int32_t>((u >> 23) & 0xFF) - 127 + 15;
+  uint32_t mant = u & 0x7FFFFF;
+  if (((u >> 23) & 0xFF) == 0xFF)
+    return static_cast<uint16_t>(sign | 0x7C00 | (mant ? 0x200 | (mant >> 13) : 0));
+  if (exp >= 0x1F) return static_cast<uint16_t>(sign | 0x7C00);
+  if (exp <= 0) {
+    if (exp < -10) return sign;
+    mant |= 0x800000;
+    uint32_t shift = static_cast<uint32_t>(14 - exp);
+    uint32_t half = mant >> shift;
+    uint32_t rem = mant & ((1u << shift) - 1);
+    uint32_t mid = 1u << (shift - 1);
+    if (rem > mid || (rem == mid && (half & 1))) ++half;
+    return static_cast<uint16_t>(sign | half);
+  }
+  uint32_t half = static_cast<uint32_t>(exp << 10) | (mant >> 13);
+  uint32_t rem = mant & 0x1FFF;
+  if (rem > 0x1000 || (rem == 0x1000 && (half & 1))) ++half;
+  return static_cast<uint16_t>(sign | half);
+}
+__device__ inline int DNearestInt(float fval) {
+  float val = fval + 12582912.0f;
+  int i = __float_as_int(val);
+  return (i & 0x007fffff) - 0x00400000;
+}
+__device__ inline float DLoadAct(const void* base, ActDT dt, int64_t idx) {
+  switch (dt) {
+    case ActDT::kF32: return static_cast<const float*>(base)[idx];
+    case ActDT::kF16: return DF16ToF32(static_cast<const uint16_t*>(base)[idx]);
+    default: return DBF16ToF32(static_cast<const uint16_t*>(base)[idx]);
+  }
+}
+
+// Q8_K (thread-per-256-superblock): cuda_quant_dot.cu QuantizeQ8KKernel.
+// The per-super-block body is factored so EVERY arm that produces Q8_K
+// activation scratch (standalone grid, MMVQ LDS prologue, norm-fused
+// epilogue) produces BYTE-IDENTICAL output: same amax first-occurrence
+// tie-break, same scale/iscale arithmetic, same bsums walk. Asserted by
+// tests/vt/test_rocm_quant_dot.cpp on random AND tied-amax inputs.
+__device__ inline void QuantQ8KSBlock(vt::cpu::BlockQ8_K& y, const void* __restrict__ a,
+                                      ActDT adt, int64_t elem0) {
+  using vt::cpu::kQK_K;
+  float mx = 0.0f, amax = 0.0f;
+  for (int j = 0; j < kQK_K; ++j) {
+    const float x = DLoadAct(a, adt, elem0 + j);
+    if (const float ax = fabsf(x); ax > amax) { amax = ax; mx = x; }
+  }
+  if (amax == 0.0f) {
+    y.d = 0.0f;
+    for (int j = 0; j < kQK_K; ++j) y.qs[j] = 0;
+    for (int g = 0; g < kQK_K / 16; ++g) y.bsums[g] = 0;
+    return;
+  }
+  const float iscale = -127.0f / mx;
+  for (int j = 0; j < kQK_K; ++j) {
+    const int v = DNearestInt(iscale * DLoadAct(a, adt, elem0 + j));
+    y.qs[j] = static_cast<int8_t>(v < 127 ? v : 127);
+  }
+  for (int g = 0; g < kQK_K / 16; ++g) {
+    int sum = 0;
+    for (int ii = 0; ii < 16; ++ii) sum += y.qs[g * 16 + ii];
+    y.bsums[g] = static_cast<int16_t>(sum);
+  }
+  y.d = 1.0f / iscale;
+}
+
+}  // namespace vt::rocm
+
+#endif  // VLLM_CPP_SRC_VT_ROCM_ROCM_ACT_QUANT_H_