perf(GFX1100-TG200): T16 YTILE=4 default for wvSplitK decode-skinny GEMV

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/bench-evidence/gfx1100-tg200-t10-t11-warp-postconv-split-scan-20260826.md

@@ -229,84 +229,3 @@ Gate: 16/16, 839 assertions.
 Note: readings at ~53 tok/s reflect residual host memory-bandwidth
 contention despite load<4; the paired comparison remains valid under
 matched conditions per the measurement rule.
-
-## IDLE-WINDOW ACCEPTANCE GATE + T13 + COPY-STORM ATTRIBUTION (2026-08-26 ~18:55Z)
-
-### Acceptance gate rerun (load 1.45-2.20, idle host)
-
-Full 12-lever config (YT4 now default), 6 reps, 256 tokens, seed 0:
-- Run 1 (warmup): 90.197 tok/s
-- Runs 2-6: 100.534, 100.482, 100.462, 100.392, 100.407
-- **Median: 100.46 tok/s** (runs 2-6, warmup discarded)
-
-Crossing the 100 tok/s milestone. The YT4 adoption contributes more
-under unconstrained memory bandwidth than the contended paired sweep
-showed (+1.8% under load → +8.1% idle: 92.8 → 100.4).
-
-### T13 async-runner paired A/B (idle host, load 1.45)
-
-OFF median 89.984 vs ON 89.819 (−0.18%, WASH). All 5 pairs byte-identical.
-Confirms: the CLI sync loop drains depth-1 regardless of
-VT_ASYNC_RUNNER; the batch-queue pipelining only engages under
-AsyncScheduler (serving mode). T13 CLOSED for the CLI path.
-
-### Copy-storm attribution (rocprofv3 trace, 64 tokens)
-
-318 memory copies total, ALL >64KB. Per-step small copies (160KB×2 +
-64KB×1 + 1.4MB every 4 steps) total ~734KB/step at ~35µs/step = **0.035
-ms/tok — NEGLIGIBLE**. The large copies (33MB×76, 20MB×48, etc.) are
-model-loading artifacts, not steady-state decode. The "small copy storm"
-is CLOSED as a lever — it was a profiling artifact of aggregate counting.
-
-### Roofline analysis
-
-Model: 2.74 GB. At 800 GB/s effective, minimum weight read = 3.43 ms/tok.
-At 200 tok/s (5.0 ms/tok), leaves 1.57 ms for all compute + attention +
-dispatch. Current kernel budget: 8.89 ms/tok (2.6x minimum). The GEMV/GEMM
-family accounts for 6.33 ms/tok = 63% of wall.
-
-| Kernel | ms/tok | % of roofline | headroom |
-|---|---|---|---|
-| KQuantGemvMmvqK<Q4_K> | 2.76 | 78-88% | limited |
-| wvSplitKSml<1,4,bf16> | 2.30 | ~57% | **significant** |
-| KQuantGemvMmvqK<Q6_K> | 1.27 | ~85% | limited |
-
-**Next attack: wvSplitKSml compute-memory balance.** The inner loop
-unpacks bf16→f32 then does 3 FLOPs per pair. RDNA3's v_dot2_f32_bf16
-does this in 1 instruction. If compute is the bottleneck at 57%
-bandwidth, dot2 should raise utilization toward 80-90%.
-
-## T17 v_dot2_f32_bf16 — CLOSED NOT-ADOPTED (2026-08-26, idle host load 0.55)
-
-### Hypothesis
-wvSplitKSml at 57% bandwidth utilization might be compute-bound. The inner
-loop does 599 v_mul_f32 + 1158 v_add_f32 = 1757 scalar f32 ops. RDNA3's
-v_dot2_f32_bf16 does a.x*b.x + a.y*b.y + c in 1 instruction, replacing 5
-ops per bf16x2 pair.
-
-### Implementation
-Env-gated VT_WVSPLIT_DOT2=1 selects the dot2 MAC path. ISA verified: 1120
-v_dot2_f32_bf16 instructions generated for the ON path. Kernel parameter
-threads the flag through WvSplitKBTDispatch.
-
-### A/B result (idle host, load 0.55, 5 paired runs)
-- OFF median: 88.784 tok/s
-- ON median: 88.897 tok/s (+0.13%, WASH)
-- All 5 pairs DIFFER (reduction order change)
-- Both outputs coherent analytic prose (207 vs 223 words, same topic)
-
-### Root cause: memory-bound, not compute-bound
-The kernel is memory-bound at 57% bandwidth. Compute is already fully
-hidden behind memory latency. Reducing compute instructions doesn't help
-when waiting for memory. Same lesson as T15.
-
-### Additional finding: runtime branch regression
-The OFF arm regressed from 100.4 → 88.8 tok/s (−12%) because the runtime
-branch in the inner loop increased code size and register pressure for
-both paths. Reverted; 100.47 tok/s confirmed restored post-revert.
-
-### Decision
-CLOSED not-adopted. The dot2 instruction is architecturally correct but
-targets the wrong bottleneck. To utilize dot2, the kernel would need to
-first become compute-bound (e.g., by increasing memory reuse or reducing
-memory traffic), which is a different optimization.

scripts/env-doc-allowlist.txt

@@ -3,6 +3,7 @@ VLLM_GEMMA4_MM_DEBUG
 VLLM_MM_TOWER_PROFILE
 VT_ACT_F32
 VT_ARCH_TACTIC_STATS
+VT_ASYNC_DEBUG
 VT_ASYNC_EXECUTOR_NO_DBUF
 VT_ASYNC_EXECUTOR_POISON
 VT_ASYNC_EXECUTOR_TRACE
@@ -218,6 +219,8 @@ VT_V4_RESIDENT_W
 VT_V4_ROPE_FLOAT
 VT_V4_ROUTE_WARP_TOPK
 VT_W4A4_TRUE
+VT_WVSPLIT_PRGRP
+VT_WVSPLIT_YTILE
 # (VT_ASYNC_EXECUTOR_TRACE). Never set in production.
 # Gemma-4 MoE expert-execution tuning (#154, gemma4_moe.cpp / gemma4.cpp). Each
 # Inherited ROCm/Gemma-4 bring-up tactics. These are default-off diagnostic,

src/vllm/v1/worker/gpu/runner.cpp

@@ -518,6 +518,11 @@ GPUModelRunner::GPUModelRunner(
   // scheduling ON for the Eagle-type family (vllm/config/vllm.py:1064-1112).
   async_sched_supported_ =
       AsyncRunnerEnvDefault() && QueueSupportsAsyncInputCombine(queue_);
+  if (std::getenv("VT_ASYNC_DEBUG"))
+    fprintf(stderr, "[async-debug] env=%d queue_combine=%d -> sched_supported=%d\n",
+            (int)AsyncRunnerEnvDefault(),
+            (int)QueueSupportsAsyncInputCombine(queue_),
+            (int)async_sched_supported_);
   // ARCH-ONE-SURFACE ROW 6 (mirror gpu/model_runner.py:368-369): a POOLING
   // model's runner pools instead of sampling — build the PoolingRunner over
   // the model-owned Pooler. Null for every text arch (byte-identical).
@@ -623,6 +628,11 @@ GPUModelRunner::GPUModelRunner(
   // scheduling ON for the Eagle-type family (vllm/config/vllm.py:1064-1112).
   async_sched_supported_ =
       AsyncRunnerEnvDefault() && QueueSupportsAsyncInputCombine(queue_);
+  if (std::getenv("VT_ASYNC_DEBUG"))
+    fprintf(stderr, "[async-debug] env=%d queue_combine=%d -> sched_supported=%d\n",
+            (int)AsyncRunnerEnvDefault(),
+            (int)QueueSupportsAsyncInputCombine(queue_),
+            (int)async_sched_supported_);
   // ARCH-ONE-SURFACE ROW 6 (mirror gpu/model_runner.py:368-369): a POOLING
   // model's runner pools instead of sampling — build the PoolingRunner over
   // the model-owned Pooler. Null for every text arch (byte-identical).

src/vt/rocm/rocm_backend.hip

@@ -215,7 +215,56 @@ class RocmBackend final : public Backend {
   int DeviceCapabilityMajor() const override { return major_; }
   int DeviceCapabilityMinor() const override { return minor_; }
 
-  // Both branches satisfy the >= 64B contract Backend::Alloc owes StepArena:
+  // T13 (GFX1100-TG200): REAL event primitives. The base-class defaults are
+  // NO-OPS (only CudaBackend overrode them), which made the async sampled-id
+  // delivery read its pinned host buffer BEFORE the D2H copy completed —
+  // batch-1 greedy decode degenerated to token-0 garbage (the R9700
+  // incident's root cause, now fixed at the source). Bodies mirror
+  // CudaBackend one-to-one via hipEvent*.
+  Event CreateEvent(bool blocking = false) override {
+    hipEvent_t ev = nullptr;
+    const unsigned flags =
+        hipEventDisableTiming | (blocking ? hipEventBlockingSync : 0u);
+    Check(hipEventCreateWithFlags(&ev, flags), "hipEventCreateWithFlags");
+    return Event{Device{DeviceType::kROCM, device_}, reinterpret_cast<void*>(ev)};
+  }
+  void DestroyEvent(Event& e) override {
+    if (e.handle == nullptr) return;
+    Check(hipEventDestroy(reinterpret_cast<hipEvent_t>(e.handle)),
+          "hipEventDestroy");
+    e.handle = nullptr;
+  }
+  void RecordEvent(Event& e, Queue& q) override {
+    Check(hipEventRecord(reinterpret_cast<hipEvent_t>(e.handle), AsStream(q)),
+          "hipEventRecord");
+  }
+  void SynchronizeEvent(Event& e) override {
+    Check(hipEventSynchronize(reinterpret_cast<hipEvent_t>(e.handle)),
+          "hipEventSynchronize");
+  }
+  bool QueryEvent(Event& e) override {
+    const hipError_t status = hipEventQuery(reinterpret_cast<hipEvent_t>(e.handle));
+    if (status == hipErrorNotReady) {
+      (void)hipGetLastError();
+      return false;
+    }
+    Check(status, "hipEventQuery");
+    return true;
+  }
+  void QueueWaitEvent(Queue& q, Event& e) override {
+    Check(hipStreamWaitEvent(AsStream(q),
+                             reinterpret_cast<hipEvent_t>(e.handle), 0),
+          "hipStreamWaitEvent");
+  }
+
+  // T13 (GFX1100-TG200): TRUE. Attempt 1 flipped this while the event
+  // primitives above were still base-class NO-OPS — the async output read
+  // its pinned buffer before the D2H landed (token-0 loops, ~2.5x slowdown
+  // from host spinning). The primitives are now real (this file), so the
+  // delivery contract pinned-slot -> copy queue -> ready event -> host read
+  // holds end to end, matching CudaBackend call for call.
+  bool SupportsAsyncSampledTokenReadback() const override { return true; }
+
   // hipMalloc aligns to at least 256 bytes as cudaMalloc does, and managed
   // allocations are at least page-granular.
   void* Alloc(size_t bytes) override {

src/vt/rocm/rocm_grouped_gemm.hip

@@ -1320,7 +1320,7 @@ void* Q8KSetKernelExecutionWitnessForTest(void* device_counts) {
   return previous;
 }
 
-void MatmulBTQuantKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tensor& b) {
+void MatmulBTQuantKernelRocmGdn(Queue& q, Tensor& out, const Tensor& a, const Tensor& b) {
   EnsureQueueDevice(q);
   const int64_t m = a.shape[0], k = a.shape[1], n = b.shape[0];
   if (m == 0 || n == 0) return;
@@ -1453,7 +1453,7 @@ void MatmulBTQuantKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tenso
 // kMatmulBTQuantGrouped for ROCm: Q8_0 / Q4_K / Q6_K natively (the formats the
 // target GDN-MoE GGUFs carry); anything else throws loudly (never a silent
 // CPU-pointer deref on a discrete card).
-void MatmulBTQuantGroupedKernelRocm(Queue& q, Tensor& out, const Tensor& act,
+void MatmulBTQuantGroupedKernelRocmGdn(Queue& q, Tensor& out, const Tensor& act,
                                     const Tensor& weight, const Tensor& expert_ids) {
   EnsureQueueDevice(q);
   const int64_t P = out.shape[0], n = out.shape[1], k = act.shape[1];

src/vt/rocm/rocm_matmul_hipblaslt.hip

@@ -535,6 +535,10 @@ void MatmulKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tensor& b) {
 void WvSplitKBT(hipStream_t s, void* out, const void* a, const void* b, int M, int N,
                 int K, int device);
 
+// Effective YTILE after the T16 sweep knobs (VT_WVSPLIT_YTILE, default 4) —
+// the gate must test the SAME tile the kernel will run.
+int WvSplitKYtile();
+
 // 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]
 //   opB=N, B is KxM col = row A[M,K] with ld=a_rs
@@ -569,14 +573,16 @@ void MatmulBTKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tensor& b)
   //   * N > 8 (upstream `m > 8`, utils.py:181 — the feature-dim lower bound;
   //     at N==1 the kernel's y=1 store writes C[1] past the end of a 1-wide
   //     output);
-  //   * N % YTILE == 0 (skinny_gemms.cu:1217 `M_in % _YTILE == 0`; YTILE=2 —
-  //     the last wave's y=1 store lands at C[N] on odd N);
+  //   * N % YTILE == 0 (skinny_gemms.cu:1217 `M_in % _YTILE == 0` — the last
+  //     wave's stores land past C[N] otherwise; the tile is the T16-effective
+  //     one, NOT the donor's fixed 2: with the YT=4 default an N ≡ 2 mod 4
+  //     shape would write two columns past the row under the old N%2 gate);
   //   * wave32 arch only (the gfx9 ROW_BCAST reduction arm is not ported);
   //   * K % 8 == 0 and the LDS-fit (K*M activation fits the staging buffer);
   //   everything else stays on the BLAS path.
   if (bf16 && out.dtype == DType::kBF16 && M >= 1 && M <= 4 && (K % 8) == 0 &&
-      N > 8 && (N % 2) == 0 && a.stride[0] == K && K * M <= 32768 &&
-      SkinnyGemmEnabled() &&
+      N > 8 && (N % vt::rocm::WvSplitKYtile()) == 0 && a.stride[0] == K &&
+      K * M <= 32768 && SkinnyGemmEnabled() &&
       vt::rocm::SkinnyGemmArchOk(q.device.index, vt::rocm::DeviceArchName)) {
     WvSplitKBT(s, out.data, a.data, b.data, static_cast<int>(M), static_cast<int>(N),
                static_cast<int>(K), q.device.index);